diff --git a/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp b/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp index a39af6ca479..660b588140f 100644 --- a/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp +++ b/cpp/include/cugraph/mtmg/detail/per_device_edgelist.hpp @@ -16,8 +16,8 @@ #pragma once -#include #include +#include // FIXME: Could use std::span once compiler supports C++20 #include @@ -301,38 +301,35 @@ class per_device_edgelist_t { if (edge_end_time_) resize_and_copy_buffers(*edge_end_time_, total_size, stream); } - auto tmp_wgt = wgt_ ? std::make_optional(std::move((*wgt_)[0])) : std::nullopt; - auto tmp_edge_id = edge_id_ ? std::make_optional(std::move((*edge_id_)[0])) : std::nullopt; - auto tmp_edge_type = - edge_type_ ? std::make_optional(std::move((*edge_type_)[0])) : std::nullopt; - auto tmp_edge_start_time = - edge_start_time_ ? std::make_optional(std::move((*edge_start_time_)[0])) : std::nullopt; - auto tmp_edge_end_time = - edge_end_time_ ? std::make_optional(std::move((*edge_end_time_)[0])) : std::nullopt; - - std::tie(store_transposed ? dst_[0] : src_[0], - store_transposed ? src_[0] : dst_[0], - tmp_wgt, - tmp_edge_id, - tmp_edge_type, - tmp_edge_start_time, - tmp_edge_end_time, - std::ignore) = - cugraph::detail::shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - handle.raft_handle(), - store_transposed ? std::move(dst_[0]) : std::move(src_[0]), - store_transposed ? std::move(src_[0]) : std::move(dst_[0]), - std::move(tmp_wgt), - std::move(tmp_edge_id), - std::move(tmp_edge_type), - std::move(tmp_edge_start_time), - std::move(tmp_edge_end_time)); - - if (tmp_wgt) ((*wgt_)[0]) = std::move(*tmp_wgt); - if (tmp_edge_id) ((*edge_id_)[0]) = std::move(*tmp_edge_id); - if (tmp_edge_type) ((*edge_type_)[0]) = std::move(*tmp_edge_type); - if (tmp_edge_start_time) ((*edge_start_time_)[0]) = std::move(*tmp_edge_start_time); - if (tmp_edge_end_time) ((*edge_end_time_)[0]) = std::move(*tmp_edge_end_time); + std::vector tmp_edge_properties{}; + if (wgt_) tmp_edge_properties.push_back(std::move((*wgt_)[0])); + if (edge_id_) tmp_edge_properties.push_back(std::move((*edge_id_)[0])); + if (edge_type_) tmp_edge_properties.push_back(std::move((*edge_type_)[0])); + if (edge_start_time_) tmp_edge_properties.push_back(std::move((*edge_start_time_)[0])); + if (edge_end_time_) tmp_edge_properties.push_back(std::move((*edge_end_time_)[0])); + + std::tie(src_[0], dst_[0], tmp_edge_properties, std::ignore) = + cugraph::shuffle_ext_edges(handle.raft_handle(), + std::move(src_[0]), + std::move(dst_[0]), + std::move(tmp_edge_properties), + store_transposed); + + size_t pos{0}; + if (wgt_) + ((*wgt_)[0]) = std::move(std::get>(tmp_edge_properties[pos++])); + if (edge_id_) + ((*edge_id_)[0]) = + std::move(std::get>(tmp_edge_properties[pos++])); + if (edge_type_) + ((*edge_type_)[0]) = + std::move(std::get>(tmp_edge_properties[pos++])); + if (edge_start_time_) + ((*edge_start_time_)[0]) = + std::move(std::get>(tmp_edge_properties[pos++])); + if (edge_end_time_) + ((*edge_end_time_)[0]) = + std::move(std::get>(tmp_edge_properties[pos++])); } private: diff --git a/cpp/include/cugraph/shuffle_functions.hpp b/cpp/include/cugraph/shuffle_functions.hpp index 4580d52bbe9..39021f076f5 100644 --- a/cpp/include/cugraph/shuffle_functions.hpp +++ b/cpp/include/cugraph/shuffle_functions.hpp @@ -15,7 +15,9 @@ */ #pragma once +#include #include +#include #include #include @@ -75,51 +77,56 @@ shuffle_ext_vertex_value_pairs(raft::handle_t const& handle, * @brief Shuffle external edges to the owning GPUs (by edge partitioning) * * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. - * @tparam edge_t Type of edge identifiers. Needs to be an integral type. - * @tparam weight_t Type of edge weight. Currently float and double are supported. - * @tparam edge_type_t Type of edge type. Needs to be an integral type, currently only int32_t is - * supported. - * @tparam edge_time_t Type of edge time. Needs to be an integral type. * * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param edge_srcs Vector of source vertex ids * @param edge_dsts Vector of destination vertex ids - * @param edge_weights Optional vector of edge weights - * @param edge_ids Optional vector of edge ids - * @param edge_types Optional vector of edge types - * @param edge_start_times Optional vector of edge start times - * @param edge_end_times Optional vector of edge end times + * @param edge_properties Vector of edge properties, each element is an arithmetic device vector * @param store_transposed Should be true if shuffled edges will be used with a cugraph::graph_t * object with store_tranposed = true. Should be false otherwise. - * @return Tuple of vectors storing edge sources, destinations, optional weights, - * optional edge ids, optional edge types, optional edge start times, optional edge end - * times mapped to this GPU and a vector storing the number of edges received from each GPU. + * @return Tuple of vectors storing edge sources, destinations, and edge properties */ -template +template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, + std::vector, std::vector> shuffle_ext_edges(raft::handle_t const& handle, rmm::device_uvector&& edge_srcs, rmm::device_uvector&& edge_dsts, - std::optional>&& edge_weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, + std::vector&& edge_properties, bool store_transposed, std::optional large_buffer_type = std::nullopt); +/** + * @ingroup graph_functions_cpp + * @brief Shuffle internal edges to the owning GPUs (by edge partitioning) + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param edge_srcs Vector of source vertex ids + * @param edge_dsts Vector of destination vertex ids + * @param edge_properties Vector of edge properties, each element is an arithmetic device vector + * @param store_transposed Should be true if shuffled edges will be used with a cugraph::graph_t + * object with store_tranposed = true. Should be false otherwise. + * @return Tuple of vectors storing edge sources, destinations, and edge properties + */ +template +std::tuple, + rmm::device_uvector, + std::vector, + std::vector> +shuffle_int_edges(raft::handle_t const& handle, + rmm::device_uvector&& majors, + rmm::device_uvector&& minors, + std::vector&& edge_properties, + bool store_transposed, + raft::host_span vertex_partition_range_lasts, + std::optional large_buffer_type = std::nullopt); + /** * @brief Shuffle local edge sources (already placed by edge partitioning) to the owning GPUs (by * vertex partitioning). @@ -129,7 +136,8 @@ shuffle_ext_edges(raft::handle_t const& handle, * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and * handles to various CUDA libraries) to run graph algorithms. * @param edge_srcs Vector of local edge source IDs - * @param vertex_partition_range_lasts Span of vertex partition range lasts (size = number of GPUs) + * @param vertex_partition_range_lasts Span of vertex partition range lasts (size = number of + * GPUs) * @param store_transposed Should be true if shuffled edges will be used with a cugraph::graph_t * object with store_tranposed = true. Should be false otherwise. * @return Vector of shuffled edge source vertex IDs (shuffled by vertex partitioning). diff --git a/cpp/src/c_api/allgather.cpp b/cpp/src/c_api/allgather.cpp index 376e598bb45..59065c5f724 100644 --- a/cpp/src/c_api/allgather.cpp +++ b/cpp/src/c_api/allgather.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-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. @@ -25,7 +25,6 @@ #include #include -#include #include #include diff --git a/cpp/src/c_api/decompress_to_edgelist.cpp b/cpp/src/c_api/decompress_to_edgelist.cpp index 9c0ceb70439..a4dc207c054 100644 --- a/cpp/src/c_api/decompress_to_edgelist.cpp +++ b/cpp/src/c_api/decompress_to_edgelist.cpp @@ -24,7 +24,6 @@ #include #include -#include #include #include diff --git a/cpp/src/c_api/ecg.cpp b/cpp/src/c_api/ecg.cpp index bc8cedc6553..e8996f3e7b8 100644 --- a/cpp/src/c_api/ecg.cpp +++ b/cpp/src/c_api/ecg.cpp @@ -17,7 +17,6 @@ #include #include -#include #include #include diff --git a/cpp/src/c_api/extract_vertex_list.cpp b/cpp/src/c_api/extract_vertex_list.cpp index e3df7bc7586..8b81560a8f4 100644 --- a/cpp/src/c_api/extract_vertex_list.cpp +++ b/cpp/src/c_api/extract_vertex_list.cpp @@ -24,7 +24,6 @@ #include #include -#include #include #include diff --git a/cpp/src/c_api/graph_functions.cpp b/cpp/src/c_api/graph_functions.cpp index 81665837aaa..2b8331b9499 100644 --- a/cpp/src/c_api/graph_functions.cpp +++ b/cpp/src/c_api/graph_functions.cpp @@ -25,7 +25,6 @@ #include #include -#include #include #include #include @@ -73,27 +72,14 @@ struct create_vertex_pairs_functor : public cugraph::c_api::abstract_functor { second_copy.data(), second_->as_type(), second_->size_, handle_.get_stream()); if constexpr (multi_gpu) { - std::tie(first_copy, - second_copy, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - cugraph::detail::shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning< - vertex_t, - edge_t, - weight_t, - edge_type_type_t, - int32_t>(handle_, - std::move(first_copy), - std::move(second_copy), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt); + std::vector edge_properties{}; + + std::tie(first_copy, second_copy, std::ignore, std::ignore) = + cugraph::shuffle_ext_edges(handle_, + std::move(first_copy), + std::move(second_copy), + std::move(edge_properties), + false); } // FIXME: use std::tuple (template) instead. result_ = new cugraph::c_api::cugraph_vertex_pairs_t{ diff --git a/cpp/src/c_api/graph_mg.cpp b/cpp/src/c_api/graph_mg.cpp index 0e495c41840..add5c0df5b7 100644 --- a/cpp/src/c_api/graph_mg.cpp +++ b/cpp/src/c_api/graph_mg.cpp @@ -146,39 +146,40 @@ struct create_graph_functor : public cugraph::c_api::abstract_functor { rmm::device_uvector edgelist_dsts = concatenate(handle_, dst_, num_arrays_); - std::optional> edgelist_weights = - weights_ ? std::make_optional(concatenate(handle_, weights_, num_arrays_)) + std::vector edgelist_edge_properties{}; + + if (weights_) + edgelist_edge_properties.push_back(concatenate(handle_, weights_, num_arrays_)); + if (edge_ids_) + edgelist_edge_properties.push_back(concatenate(handle_, edge_ids_, num_arrays_)); + if (edge_type_ids_) + edgelist_edge_properties.push_back( + concatenate(handle_, edge_type_ids_, num_arrays_)); + + std::tie(edgelist_srcs, edgelist_dsts, edgelist_edge_properties, std::ignore) = + cugraph::shuffle_ext_edges(handle_, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::move(edgelist_edge_properties), + store_transposed); + + size_t pos{0}; + auto edgelist_weights = + weights_ ? std::make_optional(std::move( + std::get>(edgelist_edge_properties[pos++]))) : std::nullopt; - - std::optional> edgelist_edge_ids = - edge_ids_ ? std::make_optional(concatenate(handle_, edge_ids_, num_arrays_)) + auto edgelist_edge_ids = + edge_ids_ ? std::make_optional(std::move( + std::get>(edgelist_edge_properties[pos++]))) : std::nullopt; - - std::optional> edgelist_edge_types = - edge_type_ids_ - ? std::make_optional(concatenate(handle_, edge_type_ids_, num_arrays_)) - : std::nullopt; + auto edgelist_edge_types = + edge_type_ids_ ? std::make_optional(std::move(std::get>( + edgelist_edge_properties[pos++]))) + : std::nullopt; std::optional> edgelist_edge_start_times{std::nullopt}; std::optional> edgelist_edge_end_times{std::nullopt}; - std::tie(edgelist_srcs, - edgelist_dsts, - edgelist_weights, - edgelist_edge_ids, - edgelist_edge_types, - edgelist_edge_start_times, - edgelist_edge_end_times, - std::ignore) = cugraph::shuffle_ext_edges(handle_, - std::move(edgelist_srcs), - std::move(edgelist_dsts), - std::move(edgelist_weights), - std::move(edgelist_edge_ids), - std::move(edgelist_edge_types), - std::move(edgelist_edge_start_times), - std::move(edgelist_edge_end_times), - store_transposed); - if (vertex_list) { vertex_list = cugraph::shuffle_ext_vertices(handle_, std::move(*vertex_list)); } diff --git a/cpp/src/c_api/hits.cpp b/cpp/src/c_api/hits.cpp index d469b18a2c5..ad813f7b983 100644 --- a/cpp/src/c_api/hits.cpp +++ b/cpp/src/c_api/hits.cpp @@ -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. @@ -18,11 +18,11 @@ #include "c_api/graph.hpp" #include "c_api/resource_handle.hpp" #include "c_api/utils.hpp" +#include "detail/shuffle_wrappers.hpp" #include #include -#include #include #include diff --git a/cpp/src/c_api/induced_subgraph.cpp b/cpp/src/c_api/induced_subgraph.cpp index 2f73652857c..e8d23792fe8 100644 --- a/cpp/src/c_api/induced_subgraph.cpp +++ b/cpp/src/c_api/induced_subgraph.cpp @@ -24,7 +24,6 @@ #include #include -#include #include #include diff --git a/cpp/src/c_api/k_core.cpp b/cpp/src/c_api/k_core.cpp index 2e967722f82..bc73335e39a 100644 --- a/cpp/src/c_api/k_core.cpp +++ b/cpp/src/c_api/k_core.cpp @@ -19,11 +19,11 @@ #include "c_api/graph.hpp" #include "c_api/resource_handle.hpp" #include "c_api/utils.hpp" +#include "detail/shuffle_wrappers.hpp" #include #include -#include #include #include diff --git a/cpp/src/c_api/k_truss.cpp b/cpp/src/c_api/k_truss.cpp index 19a11e287ed..d79ba012e9f 100644 --- a/cpp/src/c_api/k_truss.cpp +++ b/cpp/src/c_api/k_truss.cpp @@ -23,7 +23,6 @@ #include #include -#include #include #include diff --git a/cpp/src/c_api/katz.cpp b/cpp/src/c_api/katz.cpp index 545af2efaba..b99e93572d1 100644 --- a/cpp/src/c_api/katz.cpp +++ b/cpp/src/c_api/katz.cpp @@ -19,11 +19,11 @@ #include "c_api/graph.hpp" #include "c_api/resource_handle.hpp" #include "c_api/utils.hpp" +#include "detail/shuffle_wrappers.hpp" #include #include -#include #include #include diff --git a/cpp/src/c_api/leiden.cpp b/cpp/src/c_api/leiden.cpp index 662f6870b5a..60702e16a63 100644 --- a/cpp/src/c_api/leiden.cpp +++ b/cpp/src/c_api/leiden.cpp @@ -25,7 +25,6 @@ #include #include -#include #include #include diff --git a/cpp/src/c_api/lookup_src_dst.cpp b/cpp/src/c_api/lookup_src_dst.cpp index a4114dc0bd7..ec3738944a5 100644 --- a/cpp/src/c_api/lookup_src_dst.cpp +++ b/cpp/src/c_api/lookup_src_dst.cpp @@ -25,7 +25,6 @@ #include #include -#include #include #include #include diff --git a/cpp/src/c_api/louvain.cpp b/cpp/src/c_api/louvain.cpp index 156981aaa4d..1326daa4d00 100644 --- a/cpp/src/c_api/louvain.cpp +++ b/cpp/src/c_api/louvain.cpp @@ -24,7 +24,6 @@ #include #include -#include #include #include diff --git a/cpp/src/c_api/negative_sampling.cpp b/cpp/src/c_api/negative_sampling.cpp index b2944344a76..c2655066c1a 100644 --- a/cpp/src/c_api/negative_sampling.cpp +++ b/cpp/src/c_api/negative_sampling.cpp @@ -20,11 +20,11 @@ #include "c_api/random.hpp" #include "c_api/resource_handle.hpp" #include "c_api/utils.hpp" +#include "detail/shuffle_wrappers.hpp" #include #include -#include #include #include diff --git a/cpp/src/c_api/pagerank.cpp b/cpp/src/c_api/pagerank.cpp index 329c85a39c4..9ff11457921 100644 --- a/cpp/src/c_api/pagerank.cpp +++ b/cpp/src/c_api/pagerank.cpp @@ -19,11 +19,11 @@ #include "c_api/graph.hpp" #include "c_api/resource_handle.hpp" #include "c_api/utils.hpp" +#include "detail/shuffle_wrappers.hpp" #include #include -#include #include #include #include diff --git a/cpp/src/c_api/random.cpp b/cpp/src/c_api/random.cpp index 4b3fa7921a7..d49a631391d 100644 --- a/cpp/src/c_api/random.cpp +++ b/cpp/src/c_api/random.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-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. @@ -13,6 +13,7 @@ * See the License for the specific language governing permissions and * limitations under the License. */ + #include "c_api/random.hpp" #include "c_api/abstract_functor.hpp" @@ -23,7 +24,6 @@ #include -#include #include namespace { diff --git a/cpp/src/c_api/similarity.cpp b/cpp/src/c_api/similarity.cpp index b5151ea9b39..4f6f099d48d 100644 --- a/cpp/src/c_api/similarity.cpp +++ b/cpp/src/c_api/similarity.cpp @@ -23,7 +23,6 @@ #include #include -#include #include #include #include diff --git a/cpp/src/c_api/strongly_connected_components.cpp b/cpp/src/c_api/strongly_connected_components.cpp index a5e25ef0d00..6871aeeaafe 100644 --- a/cpp/src/c_api/strongly_connected_components.cpp +++ b/cpp/src/c_api/strongly_connected_components.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-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. @@ -23,7 +23,6 @@ #include #include -#include #include #include diff --git a/cpp/src/c_api/weakly_connected_components.cpp b/cpp/src/c_api/weakly_connected_components.cpp index 8744b671298..00b9660a974 100644 --- a/cpp/src/c_api/weakly_connected_components.cpp +++ b/cpp/src/c_api/weakly_connected_components.cpp @@ -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. @@ -23,7 +23,6 @@ #include #include -#include #include #include diff --git a/cpp/src/community/approx_weighted_matching_impl.cuh b/cpp/src/community/approx_weighted_matching_impl.cuh index cc1d9958f37..ac7a2b1b3e2 100644 --- a/cpp/src/community/approx_weighted_matching_impl.cuh +++ b/cpp/src/community/approx_weighted_matching_impl.cuh @@ -15,6 +15,7 @@ */ #pragma once +#include "detail/shuffle_wrappers.hpp" #include "prims/fill_edge_property.cuh" #include "prims/reduce_op.cuh" #include "prims/transform_e.cuh" @@ -23,7 +24,6 @@ #include "utilities/collect_comm.cuh" #include -#include #include #include diff --git a/cpp/src/community/detail/refine_impl.cuh b/cpp/src/community/detail/refine_impl.cuh index ad30d6af145..489230c3c77 100644 --- a/cpp/src/community/detail/refine_impl.cuh +++ b/cpp/src/community/detail/refine_impl.cuh @@ -17,16 +17,15 @@ #include "common_methods.hpp" #include "detail/graph_partition_utils.cuh" +#include "detail/shuffle_wrappers.hpp" #include "maximal_independent_moves.hpp" #include "prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh" #include "prims/per_v_transform_reduce_incoming_outgoing_e.cuh" #include "prims/reduce_op.cuh" -#include "prims/transform_reduce_e.cuh" #include "prims/transform_reduce_e_by_src_dst_key.cuh" #include "prims/update_edge_src_dst_property.cuh" #include "utilities/collect_comm.cuh" -#include #include #include #include @@ -645,24 +644,18 @@ refine_clustering( std::optional> coarse_edge_weights{std::nullopt}; if constexpr (multi_gpu) { - std::tie(d_srcs, - d_dsts, - d_weights, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - cugraph::shuffle_ext_edges( - handle, - std::move(d_srcs), - std::move(d_dsts), - std::move(d_weights), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - GraphViewType::is_storage_transposed); + std::vector edge_properties{}; + if (d_weights) edge_properties.push_back(std::move(*d_weights)); + + std::tie(d_srcs, d_dsts, edge_properties, std::ignore) = + cugraph::shuffle_ext_edges(handle, + std::move(d_srcs), + std::move(d_dsts), + std::move(edge_properties), + GraphViewType::is_storage_transposed); + + if (d_weights) + *d_weights = std::move(std::get>(edge_properties[0])); } std::tie(decision_graph, coarse_edge_weights, std::ignore, std::ignore, renumber_map) = diff --git a/cpp/src/community/edge_triangle_count_impl.cuh b/cpp/src/community/edge_triangle_count_impl.cuh index 8c98e193b87..881ac89ee42 100644 --- a/cpp/src/community/edge_triangle_count_impl.cuh +++ b/cpp/src/community/edge_triangle_count_impl.cuh @@ -22,9 +22,9 @@ #include "prims/per_v_pair_dst_nbr_intersection.cuh" #include "prims/transform_e.cuh" -#include #include #include +#include #include #include @@ -266,40 +266,26 @@ edge_property_t edge_triangle_count_impl( rmm::device_uvector pair_srcs(0, handle.get_stream()); rmm::device_uvector pair_dsts(0, handle.get_stream()); - std::optional> pair_count{std::nullopt}; + std::vector pair_properties{}; - std::optional> opt_increase_count = - std::make_optional(rmm::device_uvector(increase_count.size(), handle.get_stream())); + rmm::device_uvector pair_count(increase_count.size(), handle.get_stream()); - raft::copy((*opt_increase_count).begin(), - increase_count.begin(), - increase_count.size(), - handle.get_stream()); + raft::copy( + pair_count.begin(), increase_count.begin(), increase_count.size(), handle.get_stream()); + + pair_properties.push_back(std::move(pair_count)); // There are still multiple copies here but is it worth sorting and reducing again? - std::tie(pair_srcs, - pair_dsts, - std::ignore, - pair_count, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - handle, - std::move(std::get<0>(vertex_pair_buffer)), - std::move(std::get<1>(vertex_pair_buffer)), - std::nullopt, - // FIXME: Add general purpose function for shuffling vertex pairs and arbitrary attributes - std::move(opt_increase_count), - std::nullopt, - std::nullopt, - std::nullopt, - cur_graph_view.vertex_partition_range_lasts()); + std::tie(pair_srcs, pair_dsts, pair_properties, std::ignore) = + shuffle_int_edges(handle, + std::move(std::get<0>(vertex_pair_buffer)), + std::move(std::get<1>(vertex_pair_buffer)), + std::move(pair_properties), + false, + cur_graph_view.vertex_partition_range_lasts()); + + pair_count = std::move(std::get>(pair_properties[0])); + pair_properties.clear(); thrust::for_each( handle.get_thrust_policy(), @@ -309,7 +295,7 @@ edge_property_t edge_triangle_count_impl( num_triangles = num_triangles.data(), pair_srcs = pair_srcs.data(), pair_dsts = pair_dsts.data(), - pair_count = (*pair_count).data(), + pair_count = pair_count.data(), edge_first] __device__(auto idx) { auto src = pair_srcs[idx]; auto dst = pair_dsts[idx]; diff --git a/cpp/src/community/k_truss_impl.cuh b/cpp/src/community/k_truss_impl.cuh index 6fed8cb986b..896f2b280ba 100644 --- a/cpp/src/community/k_truss_impl.cuh +++ b/cpp/src/community/k_truss_impl.cuh @@ -25,8 +25,8 @@ #include #include -#include #include +#include #include #include @@ -458,28 +458,17 @@ k_truss(raft::handle_t const& handle, }); if constexpr (multi_gpu) { + std::vector edge_properties{}; + std::tie(std::get<0>(edgelist_to_update_count), std::get<1>(edgelist_to_update_count), std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - detail::shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - handle, - std::move(std::get<0>(edgelist_to_update_count)), - std::move(std::get<1>(edgelist_to_update_count)), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - cur_graph_view.vertex_partition_range_lasts()); + std::ignore) = shuffle_int_edges(handle, + std::move(std::get<0>(edgelist_to_update_count)), + std::move(std::get<1>(edgelist_to_update_count)), + std::move(edge_properties), + false, + cur_graph_view.vertex_partition_range_lasts()); } thrust::sort(handle.get_thrust_policy(), @@ -647,28 +636,15 @@ k_truss(raft::handle_t const& handle, // shuffle the edges if multi_gpu if constexpr (multi_gpu) { - std::tie(weak_edgelist_dsts, - weak_edgelist_srcs, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - detail::shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - handle, - std::move(weak_edgelist_dsts), - std::move(weak_edgelist_srcs), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - cur_graph_view.vertex_partition_range_lasts()); + std::vector edge_properties{}; + + std::tie(weak_edgelist_dsts, weak_edgelist_srcs, std::ignore, std::ignore) = + shuffle_int_edges(handle, + std::move(weak_edgelist_dsts), + std::move(weak_edgelist_srcs), + std::move(edge_properties), + false, + cur_graph_view.vertex_partition_range_lasts()); } thrust::sort( diff --git a/cpp/src/community/leiden_impl.cuh b/cpp/src/community/leiden_impl.cuh index 3adee40aa12..dc823c09654 100644 --- a/cpp/src/community/leiden_impl.cuh +++ b/cpp/src/community/leiden_impl.cuh @@ -56,7 +56,7 @@ vertex_t remove_duplicates(raft::handle_t const& handle, rmm::device_uvector -#include #include #include #include diff --git a/cpp/src/community/triangle_count_impl.cuh b/cpp/src/community/triangle_count_impl.cuh index 8541b5ba170..baaff6299a1 100644 --- a/cpp/src/community/triangle_count_impl.cuh +++ b/cpp/src/community/triangle_count_impl.cuh @@ -15,6 +15,7 @@ */ #pragma once +#include "detail/shuffle_wrappers.hpp" #include "prims/extract_transform_if_e.cuh" #include "prims/fill_edge_property.cuh" #include "prims/transform_e.cuh" @@ -22,7 +23,6 @@ #include "prims/update_edge_src_dst_property.cuh" #include -#include #include #include #include @@ -443,17 +443,9 @@ void triangle_count(raft::handle_t const& handle, extract_low_to_high_degree_edges_pred_op_t{}); if constexpr (multi_gpu) { - std::tie( - srcs, dsts, std::ignore, std::ignore, std::ignore, std::ignore, std::ignore, std::ignore) = - shuffle_ext_edges(handle, - std::move(srcs), - std::move(dsts), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - false); + std::vector edge_properties{}; + std::tie(srcs, dsts, std::ignore, std::ignore) = shuffle_ext_edges( + handle, std::move(srcs), std::move(dsts), std::move(edge_properties), false); } std::tie(modified_graph, std::ignore, std::ignore, std::ignore, renumber_map) = diff --git a/cpp/src/components/weakly_connected_components_impl.cuh b/cpp/src/components/weakly_connected_components_impl.cuh index a2ad243b0fd..93e8733a1db 100644 --- a/cpp/src/components/weakly_connected_components_impl.cuh +++ b/cpp/src/components/weakly_connected_components_impl.cuh @@ -15,7 +15,6 @@ */ #pragma once -#include "prims/extract_transform_if_e.cuh" #include "prims/fill_edge_property.cuh" #include "prims/fill_edge_src_dst_property.cuh" #include "prims/transform_e.cuh" @@ -26,7 +25,6 @@ #include "utilities/collect_comm.cuh" #include -#include #include #include #include @@ -858,24 +856,14 @@ void weakly_connected_components_impl(raft::handle_t const& handle, handle.get_thrust_policy(), input_first, input_first + num_inserts, output_first); if (GraphViewType::is_multi_gpu) { - std::tie(std::get<0>(edge_buffer), - std::get<1>(edge_buffer), - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - shuffle_ext_edges( - handle, - std::move(std::get<0>(edge_buffer)), - std::move(std::get<1>(edge_buffer)), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - GraphViewType::is_storage_transposed); + std::vector edge_properties{}; + + std::tie(std::get<0>(edge_buffer), std::get<1>(edge_buffer), std::ignore, std::ignore) = + shuffle_ext_edges(handle, + std::move(std::get<0>(edge_buffer)), + std::move(std::get<1>(edge_buffer)), + std::move(edge_properties), + GraphViewType::is_storage_transposed); auto edge_first = get_dataframe_buffer_begin(edge_buffer); auto edge_last = get_dataframe_buffer_end(edge_buffer); thrust::sort(handle.get_thrust_policy(), edge_first, edge_last); diff --git a/cpp/src/detail/collect_comm_wrapper.cuh b/cpp/src/detail/collect_comm_wrapper.cuh index b5e6a7dbcf7..71a6ab036bf 100644 --- a/cpp/src/detail/collect_comm_wrapper.cuh +++ b/cpp/src/detail/collect_comm_wrapper.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-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. @@ -16,16 +16,8 @@ #pragma once -#include "c_api/capi_helper.hpp" -#include "structure/detail/structure_utils.cuh" #include "utilities/collect_comm.cuh" -#include -#include - -#include -#include - namespace cugraph { namespace detail { diff --git a/cpp/src/detail/collect_comm_wrapper_mg_v32_e32.cu b/cpp/src/detail/collect_comm_wrapper_mg_v32_e32.cu index 89a028531c0..6ffc17d70e2 100644 --- a/cpp/src/detail/collect_comm_wrapper_mg_v32_e32.cu +++ b/cpp/src/detail/collect_comm_wrapper_mg_v32_e32.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-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. @@ -14,16 +14,7 @@ * limitations under the License. */ -#include "c_api/capi_helper.hpp" -#include "detail/collect_comm_wrapper.cuh" -#include "structure/detail/structure_utils.cuh" -#include "utilities/collect_comm.cuh" - -#include -#include - -#include -#include +#include "collect_comm_wrapper.cuh" namespace cugraph { namespace detail { diff --git a/cpp/src/detail/collect_comm_wrapper_mg_v64_e64.cu b/cpp/src/detail/collect_comm_wrapper_mg_v64_e64.cu index 2eb856779a7..9a976d76cdb 100644 --- a/cpp/src/detail/collect_comm_wrapper_mg_v64_e64.cu +++ b/cpp/src/detail/collect_comm_wrapper_mg_v64_e64.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-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. @@ -14,16 +14,7 @@ * limitations under the License. */ -#include "c_api/capi_helper.hpp" -#include "detail/collect_comm_wrapper.cuh" -#include "structure/detail/structure_utils.cuh" -#include "utilities/collect_comm.cuh" - -#include -#include - -#include -#include +#include "collect_comm_wrapper.cuh" namespace cugraph { namespace detail { diff --git a/cpp/src/detail/groupby_and_count.cuh b/cpp/src/detail/groupby_and_count.cuh index 9fbed0defae..e9ac009f2a1 100644 --- a/cpp/src/detail/groupby_and_count.cuh +++ b/cpp/src/detail/groupby_and_count.cuh @@ -17,9 +17,9 @@ #pragma once #include "detail/graph_partition_utils.cuh" +#include "detail/shuffle_wrappers.hpp" #include -#include #include #include #include diff --git a/cpp/src/detail/groupby_and_count_mg_v32_e32.cu b/cpp/src/detail/groupby_and_count_mg_v32_e32.cu index 6417d553d2d..421eb10426f 100644 --- a/cpp/src/detail/groupby_and_count_mg_v32_e32.cu +++ b/cpp/src/detail/groupby_and_count_mg_v32_e32.cu @@ -13,24 +13,8 @@ * See the License for the specific language governing permissions and * limitations under the License. */ -#include "detail/graph_partition_utils.cuh" -#include "detail/groupby_and_count.cuh" -#include -#include -#include -#include -#include - -#include - -#include -#include -#include -#include -#include - -#include +#include "groupby_and_count.cuh" namespace cugraph { namespace detail { diff --git a/cpp/src/detail/groupby_and_count_mg_v64_e64.cu b/cpp/src/detail/groupby_and_count_mg_v64_e64.cu index 5002b7a1af7..8f5739038ae 100644 --- a/cpp/src/detail/groupby_and_count_mg_v64_e64.cu +++ b/cpp/src/detail/groupby_and_count_mg_v64_e64.cu @@ -13,24 +13,7 @@ * See the License for the specific language governing permissions and * limitations under the License. */ -#include "detail/graph_partition_utils.cuh" -#include "detail/groupby_and_count.cuh" - -#include -#include -#include -#include -#include - -#include - -#include -#include -#include -#include -#include - -#include +#include "groupby_and_count.cuh" namespace cugraph { namespace detail { diff --git a/cpp/src/detail/permute_range.cuh b/cpp/src/detail/permute_range.cuh index 3edc5c4f0c8..2dbab9ebba0 100644 --- a/cpp/src/detail/permute_range.cuh +++ b/cpp/src/detail/permute_range.cuh @@ -17,7 +17,6 @@ #pragma once #include -#include #include #include #include diff --git a/cpp/include/cugraph/detail/shuffle_wrappers.hpp b/cpp/src/detail/shuffle_wrappers.hpp similarity index 71% rename from cpp/include/cugraph/detail/shuffle_wrappers.hpp rename to cpp/src/detail/shuffle_wrappers.hpp index 160f36994d3..1e64c6157db 100644 --- a/cpp/include/cugraph/detail/shuffle_wrappers.hpp +++ b/cpp/src/detail/shuffle_wrappers.hpp @@ -34,14 +34,10 @@ namespace detail { /** * @ingroup shuffle_wrappers_cpp - * @brief Shuffle external (i.e. before renumbering) vertex pairs (which can be edge end points) to - * their local GPUs based on edge partitioning. + * @brief Shuffle internal (i.e. renumbered) vertex pairs (which can be edge end points) and their + * properties to their local GPUs based on edge partitioning. * * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. - * @tparam edge_t Type of edge identifiers. Needs to be an integral type. - * @tparam weight_t Type of edge weights. Needs to be a floating point type. - * @tparam edge_type_t Type of edge type identifiers. Needs to be an integral type. - * @tparam edge_time_t The type of the edge time stamp. Needs to be an integral type. * * @param[in] handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, * and handles to various CUDA libraries) to run graph algorithms. @@ -50,94 +46,25 @@ namespace detail { * sparse 2D matrix using sources as major indices) or minor=>major (otherwise) and apply the edge * partitioning to determine the local GPU. * @param[in] minors Vector of second elements in vertex pairs. - * @param[in] weights Optional vector of vertex pair weight values. - * @param[in] edge_ids Optional vector of vertex pair edge id values. - * @param[in] edge_types Optional vector of vertex pair edge type values. - * @param[in] edge_start_times Optional vector of vertex pair edge start time values. - * @param[in] edge_end_times Optional vector of vertex pair edge end time values. - * @param[in] large_buffer_type Dictates the large buffer type to use in storing the shuffled vertex - * value pairs (if the value is std::nullopt, the default RMM per-device memory resource is used). - * - * @return Tuple of vectors storing shuffled major vertices, minor vertices and optional weights, - * edge ids and edge types - */ -template -std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - std::optional large_buffer_type = std::nullopt); - -/** - * @ingroup shuffle_wrappers_cpp - * @brief Shuffle internal (i.e. renumbered) vertex pairs (which can be edge end points) to their - * local GPUs based on edge partitioning. - * - * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. - * @tparam edge_t Type of edge identifiers. Needs to be an integral type. - * @tparam weight_t Type of edge weights. Needs to be a floating point type. - * @tparam edge_type_t Type of edge type identifiers. Needs to be an integral type. - * @tparam edge_time_t Type of edge time. Needs to be an integral type. - * - * @param[in] handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, - * and handles to various CUDA libraries) to run graph algorithms. - * @param[in] majors Vector of first elemetns in vertex pairs. To determine the local GPU of a - * (major, minor) pair, we assume there exists an edge from major=>minor (if we store edges in the - * sparse 2D matrix using sources as major indices) or minor=>major (otherwise) and apply the edge - * partitioning to determine the local GPU. - * @param[in] minors Vector of second elements in vertex pairs. - * @param[in] weights Optional vector of vertex pair weight values. - * @param[in] edge_ids Optional vector of vertex pair edge id values. - * @param[in] edge_types Optional vector of vertex pair edge type values. - * @param[in] edge_start_times Optional vector of vertex pair edge start time values. - * @param[in] edge_end_times Optional vector of vertex pair edge end time values. + * @param[in] edge_properties Vector of edge property vectors. * @param[in] vertex_partition_range_lasts Vector of each GPU's vertex partition range's last * (exclusive) vertex ID. * @param[in] large_buffer_type Dictates the large buffer type to use in storing the shuffled vertex * pairs (if the value is std::nullopt, the default RMM per-device memory resource is used). * - * @return Tuple of vectors storing shuffled major vertices, minor vertices and optional weights, - * edge ids and edge types and rx counts + * @return Tuple of device vectors storing shuffled major vertices, minor vertices, a vector of + * device vectors of properties and host vector of rx counts */ -template +template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, + std::vector, std::vector> shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( raft::handle_t const& handle, rmm::device_uvector&& majors, rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, + std::vector&& edge_properties, raft::host_span vertex_partition_range_lasts, std::optional large_buffer_type = std::nullopt); diff --git a/cpp/src/link_prediction/similarity_impl.cuh b/cpp/src/link_prediction/similarity_impl.cuh index b7e0ed73238..3724ef98b53 100644 --- a/cpp/src/link_prediction/similarity_impl.cuh +++ b/cpp/src/link_prediction/similarity_impl.cuh @@ -21,9 +21,9 @@ #include "prims/update_edge_src_dst_property.cuh" #include "utilities/error_check_utils.cuh" -#include #include #include +#include #include #include @@ -422,23 +422,15 @@ all_pairs_similarity(raft::handle_t const& handle, if constexpr (multi_gpu) { // shuffle vertex pairs auto vertex_partition_range_lasts = graph_view.vertex_partition_range_lasts(); - - std::tie( - v1, v2, std::ignore, std::ignore, std::ignore, std::ignore, std::ignore, std::ignore) = - detail::shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - handle, - std::move(v1), - std::move(v2), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - vertex_partition_range_lasts); + std::vector edge_properties{}; + + std::tie(v1, v2, std::ignore, std::ignore) = + shuffle_int_edges(handle, + std::move(v1), + std::move(v2), + std::move(edge_properties), + false, + vertex_partition_range_lasts); } auto score = @@ -612,23 +604,14 @@ all_pairs_similarity(raft::handle_t const& handle, if constexpr (multi_gpu) { // shuffle vertex pairs auto vertex_partition_range_lasts = graph_view.vertex_partition_range_lasts(); - - std::tie( - v1, v2, std::ignore, std::ignore, std::ignore, std::ignore, std::ignore, std::ignore) = - detail::shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - handle, - std::move(v1), - std::move(v2), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - vertex_partition_range_lasts); + std::vector edge_properties{}; + + std::tie(v1, v2, std::ignore, std::ignore) = shuffle_int_edges(handle, + std::move(v1), + std::move(v2), + std::move(edge_properties), + false, + vertex_partition_range_lasts); } auto score = diff --git a/cpp/src/lookup/lookup_src_dst_impl.cuh b/cpp/src/lookup/lookup_src_dst_impl.cuh index a9169186bfa..0d864058a56 100644 --- a/cpp/src/lookup/lookup_src_dst_impl.cuh +++ b/cpp/src/lookup/lookup_src_dst_impl.cuh @@ -14,6 +14,7 @@ * limitations under the License. */ #include "detail/graph_partition_utils.cuh" +#include "detail/shuffle_wrappers.hpp" #include "prims/extract_transform_e.cuh" #include "prims/kv_store.cuh" #include "utilities/collect_comm.cuh" @@ -21,7 +22,6 @@ #include #include #include -#include #include #include #include diff --git a/cpp/src/prims/per_v_transform_reduce_incoming_outgoing_e.cuh b/cpp/src/prims/per_v_transform_reduce_incoming_outgoing_e.cuh index 5ba7edec894..bad23e3f13d 100644 --- a/cpp/src/prims/per_v_transform_reduce_incoming_outgoing_e.cuh +++ b/cpp/src/prims/per_v_transform_reduce_incoming_outgoing_e.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2024, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -16,7 +16,6 @@ #pragma once #include "prims/detail/per_v_transform_reduce_e.cuh" -#include "prims/vertex_frontier.cuh" #include #include @@ -24,9 +23,6 @@ #include -#include -#include - namespace cugraph { /** diff --git a/cpp/src/sampling/detail/prepare_next_frontier_impl.cuh b/cpp/src/sampling/detail/prepare_next_frontier_impl.cuh index 4132e293b2d..919953c19bf 100644 --- a/cpp/src/sampling/detail/prepare_next_frontier_impl.cuh +++ b/cpp/src/sampling/detail/prepare_next_frontier_impl.cuh @@ -16,9 +16,9 @@ #pragma once +#include "detail/shuffle_wrappers.hpp" #include "sampling/detail/sampling_utils.hpp" -#include #include #include #include diff --git a/cpp/src/sampling/detail/shuffle_and_organize_output.cu b/cpp/src/sampling/detail/shuffle_and_organize_output.cu index f04d3ab15fc..55aca503f1d 100644 --- a/cpp/src/sampling/detail/shuffle_and_organize_output.cu +++ b/cpp/src/sampling/detail/shuffle_and_organize_output.cu @@ -17,8 +17,8 @@ #include "detail/graph_partition_utils.cuh" #include -#include #include +#include #include #include diff --git a/cpp/src/sampling/negative_sampling_impl.cuh b/cpp/src/sampling/negative_sampling_impl.cuh index 89dbad0e1be..b85a0afe8f8 100644 --- a/cpp/src/sampling/negative_sampling_impl.cuh +++ b/cpp/src/sampling/negative_sampling_impl.cuh @@ -16,16 +16,16 @@ #pragma once -#include "cugraph/detail/collect_comm_wrapper.hpp" -#include "cugraph/utilities/device_comm.hpp" #include "prims/reduce_v.cuh" #include "prims/update_edge_src_dst_property.cuh" #include "thrust/iterator/zip_iterator.h" #include "utilities/collect_comm.cuh" -#include +#include #include #include +#include +#include #include #include @@ -331,29 +331,15 @@ std::tuple, rmm::device_uvector> negativ if constexpr (multi_gpu) { auto vertex_partition_range_lasts = graph_view.vertex_partition_range_lasts(); - - std::tie(batch_srcs, - batch_dsts, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - detail::shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - handle, - std::move(batch_srcs), - std::move(batch_dsts), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - vertex_partition_range_lasts); + std::vector edge_properties{}; + + std::tie(batch_srcs, batch_dsts, std::ignore, std::ignore) = + shuffle_int_edges(handle, + std::move(batch_srcs), + std::move(batch_dsts), + std::move(edge_properties), + false, + vertex_partition_range_lasts); } if (remove_existing_edges) { diff --git a/cpp/src/sampling/neighbor_sampling_impl.hpp b/cpp/src/sampling/neighbor_sampling_impl.hpp index d286a934ff3..c8a364e6d6d 100644 --- a/cpp/src/sampling/neighbor_sampling_impl.hpp +++ b/cpp/src/sampling/neighbor_sampling_impl.hpp @@ -16,11 +16,8 @@ #pragma once -#include "prims/fill_edge_property.cuh" -#include "prims/transform_e.cuh" #include "sampling/detail/sampling_utils.hpp" -#include #include #include #include diff --git a/cpp/src/sampling/random_walks_impl.cuh b/cpp/src/sampling/random_walks_impl.cuh index 6907f340843..a1c75ede1b5 100644 --- a/cpp/src/sampling/random_walks_impl.cuh +++ b/cpp/src/sampling/random_walks_impl.cuh @@ -17,6 +17,7 @@ #pragma once #include "detail/graph_partition_utils.cuh" +#include "detail/shuffle_wrappers.hpp" #include "prims/detail/nbr_intersection.cuh" #include "prims/per_v_random_select_transform_outgoing_e.cuh" #include "prims/property_op_utils.cuh" @@ -24,7 +25,6 @@ #include "prims/vertex_frontier.cuh" #include -#include #include #include #include @@ -42,7 +42,6 @@ #include #include -#include #include #include #include diff --git a/cpp/src/structure/coarsen_graph_impl.cuh b/cpp/src/structure/coarsen_graph_impl.cuh index dadeb84d507..24ae5470167 100644 --- a/cpp/src/structure/coarsen_graph_impl.cuh +++ b/cpp/src/structure/coarsen_graph_impl.cuh @@ -19,7 +19,6 @@ #include "prims/update_edge_src_dst_property.cuh" #include -#include #include #include #include @@ -345,28 +344,19 @@ coarsen_graph(raft::handle_t const& handle, lower_triangular_only); // 1-2. globally shuffle + std::vector edgelist_edge_properties{}; + if (edgelist_weights) edgelist_edge_properties.push_back(std::move(*edgelist_weights)); - std::tie(edgelist_majors, - edgelist_minors, - edgelist_weights, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - cugraph::detail::shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning< - vertex_t, - edge_t, - weight_t, - int32_t, - int32_t>(handle, - std::move(edgelist_majors), - std::move(edgelist_minors), - std::move(edgelist_weights), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt); + std::tie(edgelist_majors, edgelist_minors, edgelist_edge_properties, std::ignore) = + cugraph::shuffle_ext_edges(handle, + std::move(edgelist_majors), + std::move(edgelist_minors), + std::move(edgelist_edge_properties), + false); + + if (edgelist_weights) + *edgelist_weights = + std::move(std::get>(edgelist_edge_properties[0])); // 1-3. groupby and coarsen again @@ -478,27 +468,22 @@ coarsen_graph(raft::handle_t const& handle, reversed_edgelist_majors.begin()))); } + std::vector reversed_edgelist_edge_properties{}; + if (reversed_edgelist_weights) + reversed_edgelist_edge_properties.push_back(std::move(*reversed_edgelist_weights)); + std::tie(reversed_edgelist_majors, reversed_edgelist_minors, - reversed_edgelist_weights, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - cugraph::detail::shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning< - vertex_t, - edge_t, - weight_t, - int32_t, - int32_t>(handle, - std::move(reversed_edgelist_majors), - std::move(reversed_edgelist_minors), - std::move(reversed_edgelist_weights), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt); + reversed_edgelist_edge_properties, + std::ignore) = cugraph::shuffle_ext_edges(handle, + std::move(reversed_edgelist_majors), + std::move(reversed_edgelist_minors), + std::move(reversed_edgelist_edge_properties), + false); + + if (reversed_edgelist_weights) + *reversed_edgelist_weights = + std::move(std::get>(reversed_edgelist_edge_properties[0])); auto output_offset = concatenated_edgelist_majors.size(); diff --git a/cpp/src/structure/create_graph_from_edgelist_impl.cuh b/cpp/src/structure/create_graph_from_edgelist_impl.cuh index 447eaf663f9..1fa56f475c3 100644 --- a/cpp/src/structure/create_graph_from_edgelist_impl.cuh +++ b/cpp/src/structure/create_graph_from_edgelist_impl.cuh @@ -17,10 +17,10 @@ #include "cugraph/edge_partition_view.hpp" #include "detail/graph_partition_utils.cuh" +#include "detail/shuffle_wrappers.hpp" #include "structure/detail/structure_utils.cuh" #include -#include #include #include #include diff --git a/cpp/src/structure/graph_impl.cuh b/cpp/src/structure/graph_impl.cuh index b2a4134afdb..25f5aba890a 100644 --- a/cpp/src/structure/graph_impl.cuh +++ b/cpp/src/structure/graph_impl.cuh @@ -19,7 +19,6 @@ #include "detail/graph_partition_utils.cuh" #include "structure/detail/structure_utils.cuh" -#include #include #include #include diff --git a/cpp/src/structure/graph_view_impl.cuh b/cpp/src/structure/graph_view_impl.cuh index c4569fb4865..84575ef9572 100644 --- a/cpp/src/structure/graph_view_impl.cuh +++ b/cpp/src/structure/graph_view_impl.cuh @@ -589,8 +589,8 @@ edge_t graph_view_thas_edge_mask()) { auto value_firsts = (*(this->edge_mask_view())).value_firsts(); auto edge_counts = (*(this->edge_mask_view())).edge_counts(); - assert(value_firsts.size() == 0); - assert(edge_counts.size() == 0); + assert(value_firsts.size() == 1); + assert(edge_counts.size() == 1); return static_cast(detail::count_set_bits(handle, value_firsts[0], edge_counts[0])); } else { return this->number_of_edges_; diff --git a/cpp/src/structure/renumber_edgelist_impl.cuh b/cpp/src/structure/renumber_edgelist_impl.cuh index 7d6465b54b4..070a6a9cfc9 100644 --- a/cpp/src/structure/renumber_edgelist_impl.cuh +++ b/cpp/src/structure/renumber_edgelist_impl.cuh @@ -18,7 +18,6 @@ #include "detail/graph_partition_utils.cuh" #include "prims/kv_store.cuh" -#include #include #include #include diff --git a/cpp/src/structure/select_random_vertices_impl.hpp b/cpp/src/structure/select_random_vertices_impl.hpp index 43036dbece3..31f86ac6d45 100644 --- a/cpp/src/structure/select_random_vertices_impl.hpp +++ b/cpp/src/structure/select_random_vertices_impl.hpp @@ -15,9 +15,8 @@ */ #pragma once -#include "detail/graph_partition_utils.cuh" +#include "detail/shuffle_wrappers.hpp" -#include #include #include #include diff --git a/cpp/src/structure/symmetrize_edgelist_impl.cuh b/cpp/src/structure/symmetrize_edgelist_impl.cuh index 76514822e71..caf6844be11 100644 --- a/cpp/src/structure/symmetrize_edgelist_impl.cuh +++ b/cpp/src/structure/symmetrize_edgelist_impl.cuh @@ -15,8 +15,8 @@ */ #pragma once -#include #include +#include #include #include @@ -707,26 +707,43 @@ symmetrize_edgelist(raft::handle_t const& handle, // 2. shuffle the (to-be-flipped) upper triangular edges if constexpr (multi_gpu) { + std::vector upper_triangular_edge_properties{}; + if (upper_triangular_weights) + upper_triangular_edge_properties.push_back(std::move(*upper_triangular_weights)); + if (upper_triangular_edge_ids) + upper_triangular_edge_properties.push_back(std::move(*upper_triangular_edge_ids)); + if (upper_triangular_edge_types) + upper_triangular_edge_properties.push_back(std::move(*upper_triangular_edge_types)); + if (upper_triangular_edge_start_times) + upper_triangular_edge_properties.push_back(std::move(*upper_triangular_edge_start_times)); + if (upper_triangular_edge_end_times) + upper_triangular_edge_properties.push_back(std::move(*upper_triangular_edge_end_times)); + std::tie(upper_triangular_minors, upper_triangular_majors, - upper_triangular_weights, - upper_triangular_edge_ids, - upper_triangular_edge_types, - upper_triangular_edge_start_times, - upper_triangular_edge_end_times, - std::ignore) = - shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - handle, - std::move(upper_triangular_minors), - std::move(upper_triangular_majors), - std::move(upper_triangular_weights), - std::move(upper_triangular_edge_ids), - std::move(upper_triangular_edge_types), - std::move(upper_triangular_edge_start_times), - std::move(upper_triangular_edge_end_times)); + upper_triangular_edge_properties, + std::ignore) = shuffle_ext_edges(handle, + std::move(upper_triangular_minors), + std::move(upper_triangular_majors), + std::move(upper_triangular_edge_properties), + false); + + size_t pos = 0; + if (upper_triangular_weights) + *upper_triangular_weights = + std::move(std::get>(upper_triangular_edge_properties[pos++])); + if (upper_triangular_edge_ids) + *upper_triangular_edge_ids = + std::move(std::get>(upper_triangular_edge_properties[pos++])); + if (upper_triangular_edge_types) + *upper_triangular_edge_types = std::move( + std::get>(upper_triangular_edge_properties[pos++])); + if (upper_triangular_edge_start_times) + *upper_triangular_edge_start_times = std::move( + std::get>(upper_triangular_edge_properties[pos++])); + if (upper_triangular_edge_end_times) + *upper_triangular_edge_end_times = std::move( + std::get>(upper_triangular_edge_properties[pos++])); } // 3. merge the lower triangular and the (flipped) upper triangular edges @@ -980,27 +997,43 @@ symmetrize_edgelist(raft::handle_t const& handle, } if constexpr (multi_gpu) { + std::vector upper_triangular_edge_properties{}; + if (upper_triangular_weights) + upper_triangular_edge_properties.push_back(std::move(*upper_triangular_weights)); + if (upper_triangular_edge_ids) + upper_triangular_edge_properties.push_back(std::move(*upper_triangular_edge_ids)); + if (upper_triangular_edge_types) + upper_triangular_edge_properties.push_back(std::move(*upper_triangular_edge_types)); + if (upper_triangular_edge_start_times) + upper_triangular_edge_properties.push_back(std::move(*upper_triangular_edge_start_times)); + if (upper_triangular_edge_end_times) + upper_triangular_edge_properties.push_back(std::move(*upper_triangular_edge_end_times)); + std::tie(upper_triangular_majors, upper_triangular_minors, - upper_triangular_weights, - upper_triangular_edge_ids, - upper_triangular_edge_types, - upper_triangular_edge_start_times, - upper_triangular_edge_end_times, - std::ignore) = - shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - handle, - std::move(upper_triangular_majors), - std::move(upper_triangular_minors), - std::move(upper_triangular_weights), - std::move(upper_triangular_edge_ids), - std::move(upper_triangular_edge_types), - std::move(upper_triangular_edge_start_times), - std::move(upper_triangular_edge_end_times)); + upper_triangular_edge_properties, + std::ignore) = shuffle_ext_edges(handle, + std::move(upper_triangular_majors), + std::move(upper_triangular_minors), + std::move(upper_triangular_edge_properties), + false); + + size_t pos = 0; + if (upper_triangular_weights) + *upper_triangular_weights = + std::move(std::get>(upper_triangular_edge_properties[pos++])); + if (upper_triangular_edge_ids) + *upper_triangular_edge_ids = + std::move(std::get>(upper_triangular_edge_properties[pos++])); + if (upper_triangular_edge_types) + *upper_triangular_edge_types = std::move( + std::get>(upper_triangular_edge_properties[pos++])); + if (upper_triangular_edge_start_times) + *upper_triangular_edge_start_times = std::move( + std::get>(upper_triangular_edge_properties[pos++])); + if (upper_triangular_edge_end_times) + *upper_triangular_edge_end_times = std::move( + std::get>(upper_triangular_edge_properties[pos++])); } edgelist_majors = std::move(merged_lower_triangular_majors); diff --git a/cpp/src/structure/transpose_graph_impl.cuh b/cpp/src/structure/transpose_graph_impl.cuh index 1da4e008e6c..78f0fe4def7 100644 --- a/cpp/src/structure/transpose_graph_impl.cuh +++ b/cpp/src/structure/transpose_graph_impl.cuh @@ -16,7 +16,6 @@ #pragma once -#include #include #include #include @@ -28,7 +27,6 @@ #include -#include #include #include #include @@ -88,23 +86,19 @@ transpose_graph_impl(raft::handle_t const& handle, (*renumber_map).size())); graph = graph_t(handle); - std::tie(edgelist_dsts, - edgelist_srcs, - edgelist_weights, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - shuffle_ext_edges(handle, - std::move(edgelist_dsts), - std::move(edgelist_srcs), - std::move(edgelist_weights), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - store_transposed); + std::vector edgelist_edge_properties{}; + if (edgelist_weights) edgelist_edge_properties.push_back(std::move(*edgelist_weights)); + + std::tie(edgelist_dsts, edgelist_srcs, edgelist_edge_properties, std::ignore) = + shuffle_ext_edges(handle, + std::move(edgelist_dsts), + std::move(edgelist_srcs), + std::move(edgelist_edge_properties), + store_transposed); + + if (edgelist_weights) + *edgelist_weights = + std::move(std::get>(edgelist_edge_properties[0])); graph_t transposed_graph(handle); std::optional> transposed_edge_weights{}; diff --git a/cpp/src/structure/transpose_graph_storage_impl.cuh b/cpp/src/structure/transpose_graph_storage_impl.cuh index 513bf2f0d7f..dfbdf596048 100644 --- a/cpp/src/structure/transpose_graph_storage_impl.cuh +++ b/cpp/src/structure/transpose_graph_storage_impl.cuh @@ -16,7 +16,6 @@ #pragma once -#include #include #include #include @@ -28,7 +27,6 @@ #include -#include #include #include #include @@ -88,23 +86,19 @@ transpose_graph_storage_impl(raft::handle_t const& handle, (*renumber_map).size())); graph = graph_t(handle); - std::tie(edgelist_srcs, - edgelist_dsts, - edgelist_weights, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - shuffle_ext_edges(handle, - std::move(edgelist_srcs), - std::move(edgelist_dsts), - std::move(edgelist_weights), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - !store_transposed); + std::vector edgelist_edge_properties{}; + if (edgelist_weights) edgelist_edge_properties.push_back(std::move(*edgelist_weights)); + + std::tie(edgelist_srcs, edgelist_dsts, edgelist_edge_properties, std::ignore) = + shuffle_ext_edges(handle, + std::move(edgelist_srcs), + std::move(edgelist_dsts), + std::move(edgelist_edge_properties), + !store_transposed); + + if (edgelist_weights) + *edgelist_weights = + std::move(std::get>(edgelist_edge_properties[0])); graph_t storage_transposed_graph(handle); std::optional> storage_transposed_edge_weights{}; diff --git a/cpp/src/utilities/shuffle_local_edge_srcs_dsts.cuh b/cpp/src/utilities/shuffle_local_edge_srcs_dsts.cuh index 9d540cfc8d5..88a43a665e3 100644 --- a/cpp/src/utilities/shuffle_local_edge_srcs_dsts.cuh +++ b/cpp/src/utilities/shuffle_local_edge_srcs_dsts.cuh @@ -18,7 +18,6 @@ #include "detail/graph_partition_utils.cuh" -#include #include #include diff --git a/cpp/src/utilities/shuffle_vertex_pairs.cuh b/cpp/src/utilities/shuffle_vertex_pairs.cuh index 641af290f88..26b1a2e1414 100644 --- a/cpp/src/utilities/shuffle_vertex_pairs.cuh +++ b/cpp/src/utilities/shuffle_vertex_pairs.cuh @@ -17,8 +17,8 @@ #pragma once #include "detail/graph_partition_utils.cuh" +#include "detail/shuffle_wrappers.hpp" -#include #include #include #include @@ -36,29 +36,27 @@ namespace cugraph { namespace { -template +template +struct vertex_pair_groupby_functor_t { + func_t func_; + + template + auto __device__ operator()(TupleType tup) const + { + return func_(thrust::get<0>(tup), thrust::get<1>(tup)); + } +}; + +template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, + std::vector, std::vector> shuffle_vertex_pairs_with_values_by_gpu_id_impl( raft::handle_t const& handle, rmm::device_uvector&& majors, rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, + std::vector&& edge_properties, func_t func, std::optional large_buffer_type = std::nullopt) { @@ -68,33 +66,15 @@ shuffle_vertex_pairs_with_values_by_gpu_id_impl( auto& comm = handle.get_comms(); auto const comm_size = comm.get_size(); - int edge_property_count = 0; - size_t element_size = sizeof(vertex_t) * 2; - - if (weights) { - ++edge_property_count; - element_size += sizeof(weight_t); - } + size_t element_size = sizeof(vertex_t) * 2; - if (edge_ids) { - ++edge_property_count; - element_size += sizeof(edge_t); - } - if (edge_types) { - ++edge_property_count; - element_size += sizeof(edge_type_t); - } - if (edge_start_times) { - ++edge_property_count; - element_size += sizeof(edge_time_t); - } - if (edge_end_times) { - ++edge_property_count; - element_size += sizeof(edge_time_t); + if (edge_properties.size() == 1) { + element_size += + cugraph::variant_type_dispatch(edge_properties[0], cugraph::sizeof_arithmetic_element{}); + } else if (edge_properties.size() > 1) { + element_size += sizeof(size_t); } - if (edge_property_count > 1) { element_size = sizeof(vertex_t) * 2 + sizeof(size_t); } - auto total_global_mem = handle.get_device_properties().totalGlobalMem; auto constexpr mem_frugal_ratio = 0.05; // if the expected temporary buffer size exceeds the mem_frugal_ratio of the @@ -119,154 +99,67 @@ shuffle_vertex_pairs_with_values_by_gpu_id_impl( // than the requested size). rmm::device_uvector d_tx_value_counts(0, handle.get_stream()); - - if (edge_property_count == 0) { - d_tx_value_counts = cugraph::groupby_and_count( - thrust::make_zip_iterator(majors.begin(), minors.begin()), - thrust::make_zip_iterator(majors.end(), minors.end()), - [func] __device__(auto val) { return func(thrust::get<0>(val), thrust::get<1>(val)); }, - comm_size, - mem_frugal_threshold, - handle.get_stream(), - large_buffer_type); - } else if (edge_property_count == 1) { - if (weights) { - d_tx_value_counts = cugraph::groupby_and_count( - thrust::make_zip_iterator(majors.begin(), minors.begin(), weights->begin()), - thrust::make_zip_iterator(majors.end(), minors.end(), weights->end()), - [func] __device__(auto val) { return func(thrust::get<0>(val), thrust::get<1>(val)); }, - comm_size, - mem_frugal_threshold, - handle.get_stream(), - large_buffer_type); - } else if (edge_ids) { - d_tx_value_counts = cugraph::groupby_and_count( - thrust::make_zip_iterator(majors.begin(), minors.begin(), edge_ids->begin()), - thrust::make_zip_iterator(majors.end(), minors.end(), edge_ids->end()), - [func] __device__(auto val) { return func(thrust::get<0>(val), thrust::get<1>(val)); }, - comm_size, - mem_frugal_threshold, - handle.get_stream(), - large_buffer_type); - } else if (edge_types) { - d_tx_value_counts = cugraph::groupby_and_count( - thrust::make_zip_iterator(majors.begin(), minors.begin(), edge_types->begin()), - thrust::make_zip_iterator(majors.end(), minors.end(), edge_types->end()), - [func] __device__(auto val) { return func(thrust::get<0>(val), thrust::get<1>(val)); }, - comm_size, - mem_frugal_threshold, - handle.get_stream(), - large_buffer_type); - } else if (edge_start_times) { - d_tx_value_counts = cugraph::groupby_and_count( - thrust::make_zip_iterator(majors.begin(), minors.begin(), edge_start_times->begin()), - thrust::make_zip_iterator(majors.end(), minors.end(), edge_start_times->end()), - [func] __device__(auto val) { return func(thrust::get<0>(val), thrust::get<1>(val)); }, - comm_size, - mem_frugal_threshold, - handle.get_stream(), - large_buffer_type); - } else if (edge_end_times) { - d_tx_value_counts = cugraph::groupby_and_count( - thrust::make_zip_iterator(majors.begin(), minors.begin(), edge_end_times->begin()), - thrust::make_zip_iterator(majors.end(), minors.end(), edge_end_times->end()), - [func] __device__(auto val) { return func(thrust::get<0>(val), thrust::get<1>(val)); }, - comm_size, - mem_frugal_threshold, - handle.get_stream(), - large_buffer_type); - } + vertex_pair_groupby_functor_t groupby_functor{func}; + + if (edge_properties.size() == 0) { + d_tx_value_counts = + cugraph::groupby_and_count(thrust::make_zip_iterator(majors.begin(), minors.begin()), + thrust::make_zip_iterator(majors.end(), minors.end()), + groupby_functor, + comm_size, + mem_frugal_threshold, + handle.get_stream(), + large_buffer_type); + } else if (edge_properties.size() == 1) { + d_tx_value_counts = cugraph::variant_type_dispatch( + edge_properties[0], + [&handle, + &majors, + &minors, + &groupby_functor, + &large_buffer_type, + comm_size, + mem_frugal_threshold](auto& prop) { + return cugraph::groupby_and_count( + thrust::make_zip_iterator(majors.begin(), minors.begin(), prop.begin()), + thrust::make_zip_iterator(majors.end(), minors.end(), prop.end()), + groupby_functor, + comm_size, + mem_frugal_threshold, + handle.get_stream(), + large_buffer_type); + }); } else { - auto property_position = - large_buffer_type - ? large_buffer_manager::allocate_memory_buffer(majors.size(), handle.get_stream()) - : rmm::device_uvector(majors.size(), handle.get_stream()); + rmm::device_uvector property_position(majors.size(), handle.get_stream()); detail::sequence_fill( - handle.get_stream(), property_position.data(), property_position.size(), edge_t{0}); + handle.get_stream(), property_position.data(), property_position.size(), size_t{0}); d_tx_value_counts = cugraph::groupby_and_count( thrust::make_zip_iterator(majors.begin(), minors.begin(), property_position.begin()), thrust::make_zip_iterator(majors.end(), minors.end(), property_position.end()), - [func] __device__(auto val) { return func(thrust::get<0>(val), thrust::get<1>(val)); }, + groupby_functor, comm_size, mem_frugal_threshold, handle.get_stream(), large_buffer_type); - if (weights) { - auto tmp = large_buffer_type - ? large_buffer_manager::allocate_memory_buffer( - property_position.size(), handle.get_stream()) - : rmm::device_uvector(property_position.size(), handle.get_stream()); - - thrust::gather(handle.get_thrust_policy(), - property_position.begin(), - property_position.end(), - weights->begin(), - tmp.begin()); - - weights = std::move(tmp); - } - - if (edge_ids) { - auto tmp = large_buffer_type - ? large_buffer_manager::allocate_memory_buffer(property_position.size(), - handle.get_stream()) - : rmm::device_uvector(property_position.size(), handle.get_stream()); - - thrust::gather(handle.get_thrust_policy(), - property_position.begin(), - property_position.end(), - edge_ids->begin(), - tmp.begin()); - - edge_ids = std::move(tmp); - } - - if (edge_types) { - auto tmp = large_buffer_type ? large_buffer_manager::allocate_memory_buffer( - property_position.size(), handle.get_stream()) - : rmm::device_uvector(property_position.size(), - handle.get_stream()); - - thrust::gather(handle.get_thrust_policy(), - property_position.begin(), - property_position.end(), - edge_types->begin(), - tmp.begin()); - - edge_types = std::move(tmp); - } - - if (edge_start_times) { - auto tmp = large_buffer_type ? large_buffer_manager::allocate_memory_buffer( - property_position.size(), handle.get_stream()) - : rmm::device_uvector(property_position.size(), - handle.get_stream()); - - thrust::gather(handle.get_thrust_policy(), - property_position.begin(), - property_position.end(), - edge_start_times->begin(), - tmp.begin()); - - edge_start_times = std::move(tmp); - } - - if (edge_end_times) { - auto tmp = large_buffer_type ? large_buffer_manager::allocate_memory_buffer( - property_position.size(), handle.get_stream()) - : rmm::device_uvector(property_position.size(), - handle.get_stream()); - - thrust::gather(handle.get_thrust_policy(), - property_position.begin(), - property_position.end(), - edge_end_times->begin(), - tmp.begin()); - - edge_end_times = std::move(tmp); - } + std::for_each(edge_properties.begin(), + edge_properties.end(), + [&property_position, &handle](auto& property) { + cugraph::variant_type_dispatch( + property, [&handle, &property_position](auto& prop) { + using T = typename std::remove_reference::type; + T tmp(prop.size(), handle.get_stream()); + + thrust::gather(handle.get_thrust_policy(), + property_position.begin(), + property_position.end(), + prop.begin(), + tmp.begin()); + + prop = std::move(tmp); + }); + }); } std::vector h_tx_value_counts(d_tx_value_counts.size()); @@ -279,7 +172,7 @@ shuffle_vertex_pairs_with_values_by_gpu_id_impl( std::vector rx_counts{}; if (mem_frugal_flag || - (edge_property_count > 1)) { // trade-off potential parallelism to lower peak memory + (edge_properties.size() > 1)) { // trade-off potential parallelism to lower peak memory std::tie(majors, rx_counts) = shuffle_values( comm, majors.begin(), @@ -294,134 +187,59 @@ shuffle_vertex_pairs_with_values_by_gpu_id_impl( handle.get_stream(), large_buffer_type); - if (weights) { - std::tie(weights, rx_counts) = shuffle_values( - comm, - (*weights).begin(), - raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), - handle.get_stream(), - large_buffer_type); - } - - if (edge_ids) { - std::tie(edge_ids, rx_counts) = shuffle_values( - comm, - (*edge_ids).begin(), - raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), - handle.get_stream(), - large_buffer_type); - } - - if (edge_types) { - std::tie(edge_types, rx_counts) = shuffle_values( - comm, - (*edge_types).begin(), - raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), - handle.get_stream(), - large_buffer_type); - } - - if (edge_start_times) { - std::tie(edge_start_times, rx_counts) = shuffle_values( - comm, - (*edge_start_times).begin(), - raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), - handle.get_stream(), - large_buffer_type); - } - - if (edge_end_times) { - std::tie(edge_end_times, rx_counts) = shuffle_values( - comm, - (*edge_end_times).begin(), - raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), - handle.get_stream(), - large_buffer_type); - } + std::for_each( + edge_properties.begin(), + edge_properties.end(), + [&handle, &h_tx_value_counts, &large_buffer_type, &comm](auto& property) { + cugraph::variant_type_dispatch( + property, [&handle, &h_tx_value_counts, &large_buffer_type, &comm](auto& prop) { + std::tie(prop, std::ignore) = shuffle_values( + comm, + prop.begin(), + raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), + handle.get_stream(), + large_buffer_type); + }); + }); } else { - // There is at most one edge property set - if (weights) { - std::forward_as_tuple(std::tie(majors, minors, weights), rx_counts) = shuffle_values( - comm, - thrust::make_zip_iterator(majors.begin(), minors.begin(), weights->begin()), - raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), - handle.get_stream(), - large_buffer_type); - } else if (edge_ids) { - std::forward_as_tuple(std::tie(majors, minors, edge_ids), rx_counts) = shuffle_values( - comm, - thrust::make_zip_iterator(majors.begin(), minors.begin(), edge_ids->begin()), - raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), - handle.get_stream(), - large_buffer_type); - } else if (edge_types) { - std::forward_as_tuple(std::tie(majors, minors, edge_types), rx_counts) = shuffle_values( - comm, - thrust::make_zip_iterator(majors.begin(), minors.begin(), edge_types->begin()), - raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), - handle.get_stream(), - large_buffer_type); - } else if (edge_start_times) { - std::forward_as_tuple(std::tie(majors, minors, edge_start_times), rx_counts) = shuffle_values( - comm, - thrust::make_zip_iterator(majors.begin(), minors.begin(), edge_start_times->begin()), - raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), - handle.get_stream(), - large_buffer_type); - } else if (edge_end_times) { - std::forward_as_tuple(std::tie(majors, minors, edge_end_times), rx_counts) = shuffle_values( - comm, - thrust::make_zip_iterator(majors.begin(), minors.begin(), edge_end_times->begin()), - raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), - handle.get_stream(), - large_buffer_type); - } else { + if (edge_properties.size() == 0) { std::forward_as_tuple(std::tie(majors, minors), rx_counts) = shuffle_values( comm, thrust::make_zip_iterator(majors.begin(), minors.begin()), raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), handle.get_stream(), large_buffer_type); + } else { + cugraph::variant_type_dispatch( + edge_properties[0], + [&handle, &majors, &minors, &comm, &h_tx_value_counts, &large_buffer_type](auto& prop) { + std::forward_as_tuple(std::tie(majors, minors, prop), std::ignore) = shuffle_values( + comm, + thrust::make_zip_iterator(majors.begin(), minors.begin(), prop.begin()), + raft::host_span(h_tx_value_counts.data(), h_tx_value_counts.size()), + handle.get_stream(), + large_buffer_type); + }); } } - return std::make_tuple(std::move(majors), - std::move(minors), - std::move(weights), - std::move(edge_ids), - std::move(edge_types), - std::move(edge_start_times), - std::move(edge_end_times), - std::move(rx_counts)); + return std::make_tuple( + std::move(majors), std::move(minors), std::move(edge_properties), std::move(rx_counts)); } } // namespace -namespace detail { - -template +template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, + std::vector, std::vector> -shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - std::optional large_buffer_type) +shuffle_ext_edges(raft::handle_t const& handle, + rmm::device_uvector&& edge_srcs, + rmm::device_uvector&& edge_dsts, + std::vector&& edge_properties, + bool store_transposed, + std::optional large_buffer_type) { auto& comm = handle.get_comms(); auto const comm_size = comm.get_size(); @@ -430,44 +248,40 @@ shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( auto& minor_comm = handle.get_subcomm(cugraph::partition_manager::minor_comm_name()); auto const minor_comm_size = minor_comm.get_size(); - return shuffle_vertex_pairs_with_values_by_gpu_id_impl( - handle, - std::move(majors), - std::move(minors), - std::move(weights), - std::move(edge_ids), - std::move(edge_types), - std::move(edge_start_times), - std::move(edge_start_times), - cugraph::detail::compute_gpu_id_from_ext_edge_endpoints_t{ - comm_size, major_comm_size, minor_comm_size}, - large_buffer_type); + auto majors = store_transposed ? std::move(edge_dsts) : std::move(edge_srcs); + auto minors = store_transposed ? std::move(edge_srcs) : std::move(edge_dsts); + + std::vector rx_counts{}; + std::tie(majors, minors, edge_properties, rx_counts) = + shuffle_vertex_pairs_with_values_by_gpu_id_impl( + handle, + std::move(majors), + std::move(minors), + std::move(edge_properties), + cugraph::detail::compute_gpu_id_from_ext_edge_endpoints_t{ + comm_size, major_comm_size, minor_comm_size}, + large_buffer_type); + + edge_srcs = store_transposed ? std::move(minors) : std::move(majors); + edge_dsts = store_transposed ? std::move(majors) : std::move(minors); + + return std::make_tuple( + std::move(edge_srcs), std::move(edge_dsts), std::move(edge_properties), std::move(rx_counts)); } -template +template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, + std::vector, std::vector> -shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - raft::host_span vertex_partition_range_lasts, - std::optional large_buffer_type) +shuffle_int_edges(raft::handle_t const& handle, + rmm::device_uvector&& edge_srcs, + rmm::device_uvector&& edge_dsts, + std::vector&& edge_properties, + bool store_transposed, + raft::host_span vertex_partition_range_lasts, + std::optional large_buffer_type) + { auto& comm = handle.get_comms(); auto const comm_size = comm.get_size(); @@ -483,85 +297,29 @@ shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( vertex_partition_range_lasts.size(), handle.get_stream()); - return shuffle_vertex_pairs_with_values_by_gpu_id_impl( - handle, - std::move(majors), - std::move(minors), - std::move(weights), - std::move(edge_ids), - std::move(edge_types), - std::move(edge_start_times), - std::move(edge_start_times), - cugraph::detail::compute_gpu_id_from_int_edge_endpoints_t{ - raft::device_span(d_vertex_partition_range_lasts.data(), - d_vertex_partition_range_lasts.size()), - comm_size, - major_comm_size, - minor_comm_size}, - large_buffer_type); -} - -} // namespace detail - -template -std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_edges(raft::handle_t const& handle, - rmm::device_uvector&& edge_srcs, - rmm::device_uvector&& edge_dsts, - std::optional>&& edge_weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - bool store_transposed, - std::optional large_buffer_type) -{ - auto& comm = handle.get_comms(); - auto const comm_size = comm.get_size(); - auto majors = store_transposed ? std::move(edge_dsts) : std::move(edge_srcs); auto minors = store_transposed ? std::move(edge_srcs) : std::move(edge_dsts); + std::vector rx_counts{}; - std::tie(majors, - minors, - edge_weights, - edge_ids, - edge_types, - edge_start_times, - edge_end_times, - rx_counts) = - detail::shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( + std::tie(majors, minors, edge_properties, rx_counts) = + shuffle_vertex_pairs_with_values_by_gpu_id_impl( handle, std::move(majors), std::move(minors), - std::move(edge_weights), - std::move(edge_ids), - std::move(edge_types), - std::move(edge_start_times), - std::move(edge_end_times), + std::move(edge_properties), + cugraph::detail::compute_gpu_id_from_int_edge_endpoints_t{ + raft::device_span(d_vertex_partition_range_lasts.data(), + d_vertex_partition_range_lasts.size()), + comm_size, + major_comm_size, + minor_comm_size}, large_buffer_type); + edge_srcs = store_transposed ? std::move(minors) : std::move(majors); edge_dsts = store_transposed ? std::move(majors) : std::move(minors); - return std::make_tuple(std::move(edge_srcs), - std::move(edge_dsts), - std::move(edge_weights), - std::move(edge_ids), - std::move(edge_types), - std::move(edge_start_times), - std::move(edge_end_times), - std::move(rx_counts)); + return std::make_tuple( + std::move(edge_srcs), std::move(edge_dsts), std::move(edge_properties), std::move(rx_counts)); } } // namespace cugraph diff --git a/cpp/src/utilities/shuffle_vertex_pairs_mg_v32_e32.cu b/cpp/src/utilities/shuffle_vertex_pairs_mg_v32_e32.cu index 40df9eb9988..840f116a1a7 100644 --- a/cpp/src/utilities/shuffle_vertex_pairs_mg_v32_e32.cu +++ b/cpp/src/utilities/shuffle_vertex_pairs_mg_v32_e32.cu @@ -18,240 +18,27 @@ namespace cugraph { -namespace detail { - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - raft::host_span vertex_partition_range_lasts, - std::optional large_buffer_type); - template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - raft::host_span vertex_partition_range_lasts, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - raft::host_span vertex_partition_range_lasts, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - raft::host_span vertex_partition_range_lasts, - std::optional large_buffer_type); - -} // namespace detail - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, + std::vector, std::vector> shuffle_ext_edges(raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, + rmm::device_uvector&& edge_srcs, + rmm::device_uvector&& edge_dsts, + std::vector&& edge_properties, bool store_transposed, std::optional large_buffer_type); template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, + std::vector, std::vector> -shuffle_ext_edges(raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - bool store_transposed, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_edges(raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - bool store_transposed, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_edges(raft::handle_t const& handle, +shuffle_int_edges(raft::handle_t const& handle, rmm::device_uvector&& majors, rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, + std::vector&& edge_properties, bool store_transposed, + raft::host_span vertex_partition_range_lasts, std::optional large_buffer_type); } // namespace cugraph diff --git a/cpp/src/utilities/shuffle_vertex_pairs_mg_v64_e64.cu b/cpp/src/utilities/shuffle_vertex_pairs_mg_v64_e64.cu index 2b39a939f94..fa7f8b884f7 100644 --- a/cpp/src/utilities/shuffle_vertex_pairs_mg_v64_e64.cu +++ b/cpp/src/utilities/shuffle_vertex_pairs_mg_v64_e64.cu @@ -18,240 +18,27 @@ namespace cugraph { -namespace detail { - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - raft::host_span vertex_partition_range_lasts, - std::optional large_buffer_type); - template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - raft::host_span vertex_partition_range_lasts, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - raft::host_span vertex_partition_range_lasts, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - raft::host_span vertex_partition_range_lasts, - std::optional large_buffer_type); - -} // namespace detail - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, + std::vector, std::vector> shuffle_ext_edges(raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, + rmm::device_uvector&& edge_srcs, + rmm::device_uvector&& edge_dsts, + std::vector&& edge_properties, bool store_transposed, std::optional large_buffer_type); template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, + std::vector, std::vector> -shuffle_ext_edges(raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - bool store_transposed, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_edges(raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, - bool store_transposed, - std::optional large_buffer_type); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::vector> -shuffle_ext_edges(raft::handle_t const& handle, +shuffle_int_edges(raft::handle_t const& handle, rmm::device_uvector&& majors, rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& edge_start_times, - std::optional>&& edge_end_times, + std::vector&& edge_properties, bool store_transposed, + raft::host_span vertex_partition_range_lasts, std::optional large_buffer_type); } // namespace cugraph diff --git a/cpp/src/utilities/shuffle_vertices.cuh b/cpp/src/utilities/shuffle_vertices.cuh index 29ff79b6e5d..7e0eec2bef0 100644 --- a/cpp/src/utilities/shuffle_vertices.cuh +++ b/cpp/src/utilities/shuffle_vertices.cuh @@ -17,9 +17,9 @@ #pragma once #include "detail/graph_partition_utils.cuh" +#include "detail/shuffle_wrappers.hpp" #include -#include #include #include #include @@ -171,8 +171,25 @@ shuffle_ext_vertex_value_pairs(raft::handle_t const& handle, rmm::device_uvector&& values, std::optional large_buffer_type) { - return detail::shuffle_ext_vertex_value_pairs_to_local_gpu_by_vertex_partitioning( - handle, std::move(vertices), std::move(values), large_buffer_type); + auto const comm_size = handle.get_comms().get_size(); + auto& major_comm = handle.get_subcomm(cugraph::partition_manager::major_comm_name()); + auto const major_comm_size = major_comm.get_size(); + auto& minor_comm = handle.get_subcomm(cugraph::partition_manager::minor_comm_name()); + auto const minor_comm_size = minor_comm.get_size(); + + std::vector properties{}; + properties.push_back(std::move(values)); + + std::tie(vertices, properties) = + shuffle_keys_with_properties(handle, + std::move(vertices), + std::move(properties), + cugraph::detail::compute_gpu_id_from_ext_vertex_t{ + comm_size, major_comm_size, minor_comm_size}, + large_buffer_type); + + return std::make_tuple(std::move(vertices), + std::move(std::get>(properties[0]))); } template diff --git a/cpp/tests/c_api/mg_betweenness_centrality_test.c b/cpp/tests/c_api/mg_betweenness_centrality_test.c index 0a9f71c7b3f..2d4503fc832 100644 --- a/cpp/tests/c_api/mg_betweenness_centrality_test.c +++ b/cpp/tests/c_api/mg_betweenness_centrality_test.c @@ -250,7 +250,7 @@ int test_betweenness_centrality_specific_normalized(const cugraph_resource_handl weight_t h_wgt[] = { 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; vertex_t h_seeds[] = {0, 3}; - weight_t h_result[] = {0, 0.395833, 0.16666667, 0.08333333, 0.041666667, 0.0625}; + weight_t h_result[] = {0, 0.395833, 0.16666667, 0.166667, 0.041666667, 0.0625}; return generic_betweenness_centrality_test(handle, h_src, @@ -279,7 +279,7 @@ int test_betweenness_centrality_specific_unnormalized(const cugraph_resource_han weight_t h_wgt[] = { 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f, 0.1f, 2.1f, 1.1f, 5.1f, 3.1f, 4.1f, 7.2f, 3.2f}; vertex_t h_seeds[] = {0, 3}; - weight_t h_result[] = {0, 7.91667, 3.33333, 1.666667, 0.833333, 1.25}; + weight_t h_result[] = {0, 7.91667, 3.33333, 3.33333, 0.833333, 1.25}; return generic_betweenness_centrality_test(handle, h_src, diff --git a/cpp/tests/c_api/mg_bfs_test.c b/cpp/tests/c_api/mg_bfs_test.c index db24a180774..dd458beed6d 100644 --- a/cpp/tests/c_api/mg_bfs_test.c +++ b/cpp/tests/c_api/mg_bfs_test.c @@ -48,6 +48,8 @@ int generic_bfs_test(const cugraph_resource_handle_t* p_handle, cugraph_type_erased_device_array_t* p_sources = NULL; cugraph_type_erased_device_array_view_t* p_source_view = NULL; + if (cugraph_resource_handle_get_rank(p_handle) != 0) num_seeds = 0; + ret_code = cugraph_type_erased_device_array_create(p_handle, num_seeds, INT32, &p_sources, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "p_sources create failed."); diff --git a/cpp/tests/components/mg_mis_test.cu b/cpp/tests/components/mg_mis_test.cu index ac2fd34d10e..7985a5289a1 100644 --- a/cpp/tests/components/mg_mis_test.cu +++ b/cpp/tests/components/mg_mis_test.cu @@ -153,13 +153,8 @@ class Tests_MGMaximalIndependentSet per_v_transform_reduce_outgoing_e( *handle_, mg_graph_view, - multi_gpu ? src_inclusion_cache.view() - : cugraph::detail::edge_endpoint_property_view_t( - std::vector{inclusiong_flags.data()}, - std::vector{vertex_t{0}}), - multi_gpu ? dst_inclusion_cache.view() - : cugraph::detail::edge_endpoint_property_view_t( - inclusiong_flags.data(), vertex_t{0}), + src_inclusion_cache.view(), + dst_inclusion_cache.view(), cugraph::edge_dummy_property_t{}.view(), [] __device__(auto src, auto dst, auto src_included, auto dst_included, auto wt) { return (src == dst) ? 0 : dst_included; diff --git a/cpp/tests/components/mg_vertex_coloring_test.cu b/cpp/tests/components/mg_vertex_coloring_test.cu index 9b0b0b56f3a..6b9dfbe00d1 100644 --- a/cpp/tests/components/mg_vertex_coloring_test.cu +++ b/cpp/tests/components/mg_vertex_coloring_test.cu @@ -123,13 +123,8 @@ class Tests_MGGraphColoring per_v_transform_reduce_outgoing_e( *handle_, mg_graph_view, - multi_gpu - ? src_color_cache.view() - : cugraph::detail::edge_endpoint_property_view_t( - std::vector{d_colors.data()}, std::vector{vertex_t{0}}), - multi_gpu ? dst_color_cache.view() - : cugraph::detail::edge_endpoint_property_view_t( - d_colors.data(), vertex_t{0}), + src_color_cache.view(), + dst_color_cache.view(), cugraph::edge_dummy_property_t{}.view(), [] __device__(auto src, auto dst, auto src_color, auto dst_color, cuda::std::nullopt_t) { if ((src != dst) && (src_color == dst_color)) { @@ -161,12 +156,8 @@ class Tests_MGGraphColoring size_t nr_conflicts = cugraph::transform_reduce_e( *handle_, mg_graph_view, - multi_gpu ? src_color_cache.view() - : cugraph::detail::edge_endpoint_property_view_t( - std::vector{d_colors.begin()}, std::vector{0}), - multi_gpu ? dst_color_cache.view() - : cugraph::detail::edge_endpoint_property_view_t( - d_colors.begin(), vertex_t{0}), + src_color_cache.view(), + dst_color_cache.view(), cugraph::edge_dummy_property_t{}.view(), [renumber_map = (*mg_renumber_map).data()] __device__( auto src, auto dst, auto src_color, auto dst_color, cuda::std::nullopt_t) { diff --git a/cpp/tests/components/mis_test.cu b/cpp/tests/components/mis_test.cu index a8149c05901..0edbccfc356 100644 --- a/cpp/tests/components/mis_test.cu +++ b/cpp/tests/components/mis_test.cu @@ -127,11 +127,10 @@ class Tests_SGMaximalIndependentSet per_v_transform_reduce_outgoing_e( handle, sg_graph_view, - cugraph::detail::edge_endpoint_property_view_t( - std::vector{inclusiong_flags.data()}, - std::vector{vertex_t{0}}), - cugraph::detail::edge_endpoint_property_view_t( - inclusiong_flags.data(), vertex_t{0}), + cugraph::make_edge_src_property_view( + sg_graph_view, inclusiong_flags.data(), 1), + cugraph::make_edge_dst_property_view( + sg_graph_view, inclusiong_flags.data(), 1), cugraph::edge_dummy_property_t{}.view(), [] __device__(auto src, auto dst, auto src_included, auto dst_included, auto wt) { return (src == dst) ? 0 : dst_included; diff --git a/cpp/tests/components/vertex_coloring_test.cu b/cpp/tests/components/vertex_coloring_test.cu index 9d2b5109a14..c88ad37057b 100644 --- a/cpp/tests/components/vertex_coloring_test.cu +++ b/cpp/tests/components/vertex_coloring_test.cu @@ -105,10 +105,8 @@ class Tests_SGGraphColoring per_v_transform_reduce_outgoing_e( handle, sg_graph_view, - cugraph::detail::edge_endpoint_property_view_t( - std::vector{d_colors.data()}, std::vector{vertex_t{0}}), - cugraph::detail::edge_endpoint_property_view_t(d_colors.data(), - vertex_t{0}), + cugraph::make_edge_src_property_view(sg_graph_view, d_colors.data(), 1), + cugraph::make_edge_dst_property_view(sg_graph_view, d_colors.data(), 1), cugraph::edge_dummy_property_t{}.view(), [] __device__(auto src, auto dst, auto src_color, auto dst_color, cuda::std::nullopt_t) { if ((src != dst) && (src_color == dst_color)) { @@ -140,10 +138,8 @@ class Tests_SGGraphColoring edge_t nr_conflicts = cugraph::transform_reduce_e( handle, sg_graph_view, - cugraph::detail::edge_endpoint_property_view_t( - std::vector{d_colors.begin()}, std::vector{vertex_t{0}}), - cugraph::detail::edge_endpoint_property_view_t(d_colors.begin(), - vertex_t{0}), + cugraph::make_edge_src_property_view(sg_graph_view, d_colors.data(), 1), + cugraph::make_edge_dst_property_view(sg_graph_view, d_colors.data(), 1), cugraph::edge_dummy_property_t{}.view(), [renumber_map = (*sg_renumber_map).data()] __device__( auto src, auto dst, auto src_color, auto dst_color, cuda::std::nullopt_t) { diff --git a/cpp/tests/link_prediction/mg_similarity_test.cpp b/cpp/tests/link_prediction/mg_similarity_test.cpp index b734e787de6..6c9ef533f05 100644 --- a/cpp/tests/link_prediction/mg_similarity_test.cpp +++ b/cpp/tests/link_prediction/mg_similarity_test.cpp @@ -23,7 +23,7 @@ #include "utilities/thrust_wrapper.hpp" #include -#include +#include #include struct Similarity_Usecase { @@ -147,23 +147,14 @@ class Tests_MGSimilarity true); std::tie(v1, v2) = cugraph::test::remove_self_loops(*handle_, std::move(v1), std::move(v2)); - - std::tie( - v1, v2, std::ignore, std::ignore, std::ignore, std::ignore, std::ignore, std::ignore) = - cugraph::detail::shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning< - vertex_t, - edge_t, - weight_t, - int32_t, - int32_t>(*handle_, - std::move(v1), - std::move(v2), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - mg_graph_view.vertex_partition_range_lasts()); + std::vector edge_properties{}; + std::tie(v1, v2, std::ignore, std::ignore) = + cugraph::shuffle_int_edges(*handle_, + std::move(v1), + std::move(v2), + std::move(edge_properties), + false, + mg_graph_view.vertex_partition_range_lasts()); std::tuple, raft::device_span> vertex_pairs{ {v1.data(), v1.size()}, {v2.data(), v2.size()}}; diff --git a/cpp/tests/link_prediction/mg_weighted_similarity_test.cpp b/cpp/tests/link_prediction/mg_weighted_similarity_test.cpp index 6fa66a1a01e..4909bd9a0d4 100644 --- a/cpp/tests/link_prediction/mg_weighted_similarity_test.cpp +++ b/cpp/tests/link_prediction/mg_weighted_similarity_test.cpp @@ -23,7 +23,7 @@ #include "utilities/thrust_wrapper.hpp" #include -#include +#include #include struct Weighted_Similarity_Usecase { @@ -108,22 +108,15 @@ class Tests_MGSimilarity auto d_v1 = cugraph::test::to_device(*handle_, h_v1); auto d_v2 = std::move(two_hop_nbrs); - std::tie( - d_v1, d_v2, std::ignore, std::ignore, std::ignore, std::ignore, std::ignore, std::ignore) = - cugraph::detail::shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning< - vertex_t, - edge_t, - weight_t, - int32_t, - int32_t>(*handle_, - std::move(d_v1), - std::move(d_v2), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - mg_graph_view.vertex_partition_range_lasts()); + std::vector edge_properties{}; + + std::tie(d_v1, d_v2, std::ignore, std::ignore) = + cugraph::shuffle_int_edges(*handle_, + std::move(d_v1), + std::move(d_v2), + std::move(edge_properties), + false, + mg_graph_view.vertex_partition_range_lasts()); std::tuple, raft::device_span> vertex_pairs{ {d_v1.data(), d_v1.size()}, {d_v2.data(), d_v2.size()}}; diff --git a/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu b/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu index 644abfc5779..760a249f79c 100644 --- a/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu +++ b/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu @@ -23,11 +23,11 @@ #include "utilities/property_generator_utilities.hpp" #include "utilities/test_graphs.hpp" -#include #include #include #include #include +#include #include #include #include @@ -148,28 +148,18 @@ class Tests_MGPerVPairTransformDstNbrIntersection }); auto h_vertex_partition_range_lasts = mg_graph_view.vertex_partition_range_lasts(); + std::vector edge_properties{}; + std::tie(std::get<0>(mg_vertex_pair_buffer), std::get<1>(mg_vertex_pair_buffer), std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore, std::ignore) = - cugraph::detail::shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning< - vertex_t, - edge_t, - weight_t, - int32_t, - int32_t>(*handle_, - std::move(std::get<0>(mg_vertex_pair_buffer)), - std::move(std::get<1>(mg_vertex_pair_buffer)), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - h_vertex_partition_range_lasts); + cugraph::shuffle_int_edges(*handle_, + std::move(std::get<0>(mg_vertex_pair_buffer)), + std::move(std::get<1>(mg_vertex_pair_buffer)), + std::move(edge_properties), + false, + h_vertex_partition_range_lasts); auto mg_result_buffer = cugraph::allocate_dataframe_buffer>( cugraph::size_dataframe_buffer(mg_vertex_pair_buffer), handle_->get_stream()); diff --git a/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_weighted_intersection.cu b/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_weighted_intersection.cu index a962565a330..4969b01c78b 100644 --- a/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_weighted_intersection.cu +++ b/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_weighted_intersection.cu @@ -23,11 +23,11 @@ #include "utilities/test_graphs.hpp" #include "utilities/thrust_wrapper.hpp" -#include #include #include #include #include +#include #include #include #include @@ -174,28 +174,18 @@ class Tests_MGPerVPairTransformDstNbrIntersection }); auto h_vertex_partition_range_lasts = mg_graph_view.vertex_partition_range_lasts(); + std::vector edge_properties{}; + std::tie(std::get<0>(mg_vertex_pair_buffer), std::get<1>(mg_vertex_pair_buffer), std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore, std::ignore) = - cugraph::detail::shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning< - vertex_t, - edge_t, - weight_t, - int32_t, - int32_t>(*handle_, - std::move(std::get<0>(mg_vertex_pair_buffer)), - std::move(std::get<1>(mg_vertex_pair_buffer)), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - h_vertex_partition_range_lasts); + cugraph::shuffle_int_edges(*handle_, + std::move(std::get<0>(mg_vertex_pair_buffer)), + std::move(std::get<1>(mg_vertex_pair_buffer)), + std::move(edge_properties), + store_transposed, + h_vertex_partition_range_lasts); auto mg_result_buffer = cugraph::allocate_dataframe_buffer>( cugraph::size_dataframe_buffer(mg_vertex_pair_buffer), handle_->get_stream()); diff --git a/cpp/tests/prims/mg_transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cu b/cpp/tests/prims/mg_transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cu index 9c40be4944d..1681cdbc1e7 100644 --- a/cpp/tests/prims/mg_transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cu +++ b/cpp/tests/prims/mg_transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cu @@ -23,7 +23,6 @@ #include "utilities/property_generator_utilities.hpp" #include "utilities/test_graphs.hpp" -#include #include #include #include diff --git a/cpp/tests/sampling/mg_biased_neighbor_sampling.cpp b/cpp/tests/sampling/mg_biased_neighbor_sampling.cpp index 89f558cb5d6..b3c14a1378b 100644 --- a/cpp/tests/sampling/mg_biased_neighbor_sampling.cpp +++ b/cpp/tests/sampling/mg_biased_neighbor_sampling.cpp @@ -15,12 +15,12 @@ */ #include "detail/nbr_sampling_validate.hpp" +#include "detail/shuffle_wrappers.hpp" #include "utilities/base_fixture.hpp" #include "utilities/device_comm_wrapper.hpp" #include "utilities/mg_utilities.hpp" #include "utilities/property_generator_utilities.hpp" -#include #include #include diff --git a/cpp/tests/sampling/mg_heterogeneous_biased_neighbor_sampling.cpp b/cpp/tests/sampling/mg_heterogeneous_biased_neighbor_sampling.cpp index 5ea5cf4a15d..7682f6dccef 100644 --- a/cpp/tests/sampling/mg_heterogeneous_biased_neighbor_sampling.cpp +++ b/cpp/tests/sampling/mg_heterogeneous_biased_neighbor_sampling.cpp @@ -15,13 +15,13 @@ */ #include "detail/nbr_sampling_validate.hpp" +#include "detail/shuffle_wrappers.hpp" #include "utilities/base_fixture.hpp" #include "utilities/device_comm_wrapper.hpp" #include "utilities/mg_utilities.hpp" #include "utilities/property_generator_utilities.hpp" #include "utilities/test_graphs.hpp" -#include #include #include diff --git a/cpp/tests/sampling/mg_heterogeneous_uniform_neighbor_sampling.cpp b/cpp/tests/sampling/mg_heterogeneous_uniform_neighbor_sampling.cpp index 39dbb298d6d..4d1a3eddd0b 100644 --- a/cpp/tests/sampling/mg_heterogeneous_uniform_neighbor_sampling.cpp +++ b/cpp/tests/sampling/mg_heterogeneous_uniform_neighbor_sampling.cpp @@ -15,13 +15,13 @@ */ #include "detail/nbr_sampling_validate.hpp" +#include "detail/shuffle_wrappers.hpp" #include "utilities/base_fixture.hpp" #include "utilities/device_comm_wrapper.hpp" #include "utilities/mg_utilities.hpp" #include "utilities/property_generator_utilities.hpp" #include "utilities/test_graphs.hpp" -#include #include #include diff --git a/cpp/tests/sampling/mg_homogeneous_biased_neighbor_sampling.cpp b/cpp/tests/sampling/mg_homogeneous_biased_neighbor_sampling.cpp index 76d4156c744..f084fcc3df2 100644 --- a/cpp/tests/sampling/mg_homogeneous_biased_neighbor_sampling.cpp +++ b/cpp/tests/sampling/mg_homogeneous_biased_neighbor_sampling.cpp @@ -15,13 +15,13 @@ */ #include "detail/nbr_sampling_validate.hpp" +#include "detail/shuffle_wrappers.hpp" #include "utilities/base_fixture.hpp" #include "utilities/device_comm_wrapper.hpp" #include "utilities/mg_utilities.hpp" #include "utilities/property_generator_utilities.hpp" #include "utilities/test_graphs.hpp" -#include #include #include diff --git a/cpp/tests/sampling/mg_homogeneous_uniform_neighbor_sampling.cpp b/cpp/tests/sampling/mg_homogeneous_uniform_neighbor_sampling.cpp index e4e0a45ffd3..4135fb4136f 100644 --- a/cpp/tests/sampling/mg_homogeneous_uniform_neighbor_sampling.cpp +++ b/cpp/tests/sampling/mg_homogeneous_uniform_neighbor_sampling.cpp @@ -15,13 +15,13 @@ */ #include "detail/nbr_sampling_validate.hpp" +#include "detail/shuffle_wrappers.hpp" #include "utilities/base_fixture.hpp" #include "utilities/device_comm_wrapper.hpp" #include "utilities/mg_utilities.hpp" #include "utilities/property_generator_utilities.hpp" #include "utilities/test_graphs.hpp" -#include #include #include diff --git a/cpp/tests/sampling/mg_uniform_neighbor_sampling.cpp b/cpp/tests/sampling/mg_uniform_neighbor_sampling.cpp index 4ab198f3b10..c35acb4cc91 100644 --- a/cpp/tests/sampling/mg_uniform_neighbor_sampling.cpp +++ b/cpp/tests/sampling/mg_uniform_neighbor_sampling.cpp @@ -15,13 +15,13 @@ */ #include "detail/nbr_sampling_validate.hpp" +#include "detail/shuffle_wrappers.hpp" #include "utilities/base_fixture.hpp" #include "utilities/device_comm_wrapper.hpp" #include "utilities/mg_utilities.hpp" #include "utilities/property_generator_utilities.hpp" #include "utilities/test_graphs.hpp" -#include #include #include diff --git a/cpp/tests/structure/mg_has_edge_and_compute_multiplicity_test.cpp b/cpp/tests/structure/mg_has_edge_and_compute_multiplicity_test.cpp index 5951e0f0fb2..32dbfe9fdde 100644 --- a/cpp/tests/structure/mg_has_edge_and_compute_multiplicity_test.cpp +++ b/cpp/tests/structure/mg_has_edge_and_compute_multiplicity_test.cpp @@ -19,13 +19,12 @@ #include "utilities/device_comm_wrapper.hpp" #include "utilities/mg_utilities.hpp" #include "utilities/test_graphs.hpp" -#include "utilities/thrust_wrapper.hpp" #include -#include #include #include #include +#include #include #include @@ -119,29 +118,15 @@ class Tests_MGHasEdgeAndComputeMultiplicity vertex_t{0}, mg_graph_view.number_of_vertices(), rng_state); - - std::tie(store_transposed ? d_mg_edge_dsts : d_mg_edge_srcs, - store_transposed ? d_mg_edge_srcs : d_mg_edge_dsts, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - cugraph::detail::shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning< - vertex_t, - edge_t, - weight_t, - edge_type_t, - edge_time_t>(*handle_, - std::move(store_transposed ? d_mg_edge_dsts : d_mg_edge_srcs), - std::move(store_transposed ? d_mg_edge_srcs : d_mg_edge_dsts), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - mg_graph_view.vertex_partition_range_lasts()); + std::vector edge_properties{}; + + std::tie(d_mg_edge_srcs, d_mg_edge_dsts, std::ignore, std::ignore) = + cugraph::shuffle_int_edges(*handle_, + std::move(d_mg_edge_srcs), + std::move(d_mg_edge_dsts), + std::move(edge_properties), + store_transposed, + mg_graph_view.vertex_partition_range_lasts()); // 3. run MG has_edge & compute_multiplicity diff --git a/cpp/tests/traversal/mg_graph500_bfs_test.cu b/cpp/tests/traversal/mg_graph500_bfs_test.cu index 0de3dd1f8b7..e0f42373d50 100644 --- a/cpp/tests/traversal/mg_graph500_bfs_test.cu +++ b/cpp/tests/traversal/mg_graph500_bfs_test.cu @@ -15,6 +15,7 @@ */ #include "detail/graph_partition_utils.cuh" +#include "detail/shuffle_wrappers.hpp" #include "nbr_unrenumber_cache.cuh" #include "prims/count_if_e.cuh" #include "prims/extract_transform_if_e.cuh" @@ -32,7 +33,6 @@ #include "utilities/thrust_wrapper.hpp" #include -#include #include #include #include @@ -765,24 +765,14 @@ class Tests_GRAPH500_MGBFS std::nullopt, std::nullopt); - std::tie(src_chunks[i], - dst_chunks[i], - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - cugraph::shuffle_ext_edges( - *handle_, - std::move(src_chunks[i]), - std::move(dst_chunks[i]), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - store_transposed); + std::vector dummy_edge_property_chunk{}; + + std::tie(src_chunks[i], dst_chunks[i], dummy_edge_property_chunk, std::ignore) = + cugraph::shuffle_ext_edges(*handle_, + std::move(src_chunks[i]), + std::move(dst_chunks[i]), + std::move(dummy_edge_property_chunk), + store_transposed); } std::tie( @@ -2007,28 +1997,15 @@ class Tests_GRAPH500_MGBFS comm, num_invalids, raft::comms::op_t::SUM, handle_->get_stream()); ASSERT_EQ(num_invalids, 0) << "predecessor->v missing in the input graph."; - std::tie(query_preds, - query_vertices, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - cugraph::detail::shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning< - vertex_t, - edge_t, - weight_t, - edge_type_t, - edge_time_t>(*handle_, - std::move(query_preds), - std::move(query_vertices), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - mg_subgraph_view.vertex_partition_range_lasts()); + std::vector edge_properties{}; + + std::tie(query_preds, query_vertices, std::ignore, std::ignore) = + cugraph::shuffle_int_edges(*handle_, + std::move(query_preds), + std::move(query_vertices), + std::move(edge_properties), + store_transposed, + mg_subgraph_view.vertex_partition_range_lasts()); auto flags = mg_subgraph_view.has_edge( *handle_, diff --git a/cpp/tests/utilities/conversion_utilities_impl.cuh b/cpp/tests/utilities/conversion_utilities_impl.cuh index 2cfd9b6857d..435b133c0d3 100644 --- a/cpp/tests/utilities/conversion_utilities_impl.cuh +++ b/cpp/tests/utilities/conversion_utilities_impl.cuh @@ -20,7 +20,6 @@ #include "utilities/device_comm_wrapper.hpp" #include "utilities/thrust_wrapper.hpp" -#include #include #include #include diff --git a/cpp/tests/utilities/test_graphs.hpp b/cpp/tests/utilities/test_graphs.hpp index 86a0c63db80..4b13562a3db 100644 --- a/cpp/tests/utilities/test_graphs.hpp +++ b/cpp/tests/utilities/test_graphs.hpp @@ -452,25 +452,20 @@ class Rmat_Usecase : public detail::TranslateGraph_Usecase { } if (multi_gpu && shuffle) { - std::tie(tmp_src_v, - tmp_dst_v, - tmp_weights_v, - std::ignore, - std::ignore, - std::ignore, - std::ignore, - std::ignore) = - cugraph::shuffle_ext_edges( - handle, - std::move(tmp_src_v), - std::move(tmp_dst_v), - std::move(tmp_weights_v), - std::nullopt, - std::nullopt, - std::nullopt, - std::nullopt, - store_transposed, - large_edge_buffer_type); + std::vector tmp_edge_properties{}; + if (tmp_weights_v) tmp_edge_properties.push_back(std::move(*tmp_weights_v)); + + std::tie(tmp_src_v, tmp_dst_v, tmp_edge_properties, std::ignore) = + cugraph::shuffle_ext_edges(handle, + std::move(tmp_src_v), + std::move(tmp_dst_v), + std::move(tmp_edge_properties), + store_transposed, + large_edge_buffer_type); + + if (tmp_weights_v) + *tmp_weights_v = + std::move(std::get>(tmp_edge_properties[0])); } edge_src_chunks.push_back(std::move(tmp_src_v));