diff --git a/cpp/include/cugraph/graph_functions.hpp b/cpp/include/cugraph/graph_functions.hpp index c03d3fceb5..b45852287f 100644 --- a/cpp/include/cugraph/graph_functions.hpp +++ b/cpp/include/cugraph/graph_functions.hpp @@ -412,15 +412,58 @@ decompress_to_edgelist( /** * @ingroup graph_functions_cpp - * @brief Symmetrize edgelist. + + * @brief Identify if an edge property is part of the uniqueness of the edge, or just a value to be + combined. + */ +enum class edge_key_selector_t { + KEY, /** This property is part of the uniqueness key of the edge. */ + VALUE /** This property is a value to be combined. */ +}; + +/** + * @ingroup graph_functions_cpp + * @brief Rule for combining edge properties when merging reciprocal edges. + * + * Each edge property in the edge tuple can have a different combining rule. + * For example, the weight of an edge can be combined differently than the edge type. + * + * @param SUM Sum the values of the edge properties. + * @param MAX Take the maximum of the values of the edge properties. + * @param MIN Take the minimum of the values of the edge properties. + * @param MEAN Take the mean of the values of the edge properties. + * @param FIRST Use the value of the first edge property based on the original order of the edges in + * the input data. For an inverse edge things will be ordered based on the order of the + * original edge that the inverse edge was created from. + * @param LAST Use the value of the last edge property based on the original order of the edges in + * the input data. For an inverse edge things will be ordered based on the order of the + * original edge that the inverse edge was created from. + * @param NONE Do not combine the values, so the edge property will be unchanged, so the edge + * properties will potentially be different in each direction. If NONE is specified and there are + * multiple values in the same direction, then we will not remove any edge tuples from the edgelist + * and the resulting edge list can have duplicate edges. + */ +enum class edge_value_combining_rule_t { + SUM, /** Sum the values of the two edge properties. */ + MAX, /** Take the maximum of the values of the two edge properties. */ + MIN, /** Take the minimum of the values of the two edge properties. */ + MEAN, /** Take the mean of the values of the two edge properties. */ + FIRST, /** Use the value of the first edge property. */ + LAST, /** Use the value of the last edge property. */ + ANY, /** Use the value of any edge property. */ + NONE /** Do not combine the values. */ +}; + +/** + * @ingroup graph_functions_cpp + * @brief Add reciprocal edges and merge the edges. This was formerly called symmetrize_edgelist + * with reciprocal set to false. + * + * The input edge list is replicated as reciprocal edges. Then the original edges and the + * reciprocal edges are merged based on the provided @p edge_key_selector and @p + * edge_value_combining_rule. * * @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. - * @tparam store_transposed Flag indicating whether to use sources (if false) or destinations (if - * true) as major indices in storing edges using a 2D sparse matrix. * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) * or multi-GPU (true). * @@ -430,39 +473,74 @@ decompress_to_edgelist( * compute_gpu_id_from_ext_edge_endpoints_t to every edge should return the local GPU ID for this * function to work (edges should be pre-shuffled). * @param edgelist_dsts Vector of edge destination vertex IDs. - * @param edgelist_weights Vector of edge weights. - * @param edgelist_edge_ids Vector of edge ids - * @param edgelist_edge_types Vector of edge types - * @param edgelist_edge_start_times Vector of edge start times - * @param edgelist_edge_end_times Vector of edge end times - * @param reciprocal Flag indicating whether to keep (if set to `false`) or discard (if set to - * `true`) edges that appear only in one direction. - * @return Tuple of symmetrized sources, destinations, optional weights, optional edge ids, optional - * edge types, optional edge start times and optional edge end times + * @param edgelist_edge_properties Vector of edge properties. + * @param edge_key_selector Vector of selectors for selecting the key of the edge properties. + * @param edge_value_combining_rule Vector of rules for combining edge values when merging + * reciprocal edges. reciprocal edges. + * @param store_transposed Flag indicating whether to use sources (if false) or destinations (if + * true) as major indices in storing edges using a 2D sparse matrix. + * @return Tuple of symmetrized sources, destinations, and edge properties. */ -template +template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -symmetrize_edgelist(raft::handle_t const& handle, - rmm::device_uvector&& edgelist_srcs, - rmm::device_uvector&& edgelist_dsts, - std::optional>&& edgelist_weights, - std::optional>&& edgelist_edge_ids, - std::optional>&& edgelist_edge_types, - std::optional>&& edgelist_edge_start_times, - std::optional>&& edgelist_edge_end_times, - bool reciprocal); + std::vector> +add_reciprocals_and_merge_edge_tuples( + raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::vector&& edgelist_edge_properties, + std::vector const& edge_key_selector, + std::vector const& edge_value_combining_rule, + bool store_transposed); + +/** + * @ingroup graph_functions_cpp + * @brief Drop non-reciprocal edge tuples from the edgelist. This was formerly called + * symmetrize_edgelist with reciprocal set to true. + * + * The input edge list scanned, each edge tuple that does not have a matching reciprocal edge is + * dropped. @p edge_key_selector is used to determine if an edge property is part of the uniqueness + * key of the edge. If it is part of the uniqueness key, the edge tuple is removed if it does not + * have a matching reciprocal edge. If it is not part of the uniqueness key, the edge property is + * ignored in the comparison. + * + * @p edge_value_combining_rule is used to combine the edge values when merging reciprocal edges. + * Specifically, if the edge tuple has a matching reciprocal edge, the edge values could be + * different, so we need a rule to combine the edge values. If NONE is specified as the value + * combining rule, the edge property will be unchanged, so will potentially be different in each + * direction. Otherwise the edge property will be combined according to the specified rule so that + * both directions match. + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * or multi-GPU (true). + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param edgelist_srcs Vector of edge source vertex IDs. If multi-GPU, applying the + * compute_gpu_id_from_ext_edge_endpoints_t to every edge should return the local GPU ID for this + * function to work (edges should be pre-shuffled). + * @param edgelist_dsts Vector of edge destination vertex IDs. + * @param edgelist_edge_properties Vector of edge properties. + * @param edge_key_selector Vector of selectors for selecting the key of the edge properties. + * @param edge_value_combining_rule Vector of rules for combining edge values when merging + * reciprocal edges. reciprocal edges. + * @param store_transposed Flag indicating whether to use sources (if false) or destinations (if + * true) as major indices in storing edges using a 2D sparse matrix. + * @return Tuple of non-reciprocal sources, destinations, and edge properties. + */ +template +std::tuple, + rmm::device_uvector, + std::vector> +drop_nonreciprocal_edge_tuples( + raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::vector&& edgelist_edge_properties, + std::vector const& edge_key_selector, + std::vector const& edge_value_combining_rule, + bool store_transposed); /** * @ingroup graph_functions_cpp diff --git a/cpp/src/structure/symmetrize_edgelist_impl.cuh b/cpp/src/structure/symmetrize_edgelist_impl.cuh index e4f9dee981..0342102ebd 100644 --- a/cpp/src/structure/symmetrize_edgelist_impl.cuh +++ b/cpp/src/structure/symmetrize_edgelist_impl.cuh @@ -186,6 +186,150 @@ struct pick_from_lower_upper_t { namespace detail { +template +std::tuple, + rmm::device_uvector, + std::vector> +merge_lower_triangular( + raft::handle_t const& handle, + rmm::device_uvector&& lower_triangular_majors, + rmm::device_uvector&& lower_triangular_minors, + std::vector&& lower_triangular_properties, + rmm::device_uvector&& upper_triangular_majors, + rmm::device_uvector&& upper_triangular_minors, + std::vector&& upper_triangular_properties, + std::vector const& edge_property_combining_rules, + size_t num_lower_triangular_edges, + bool remove_edges) +{ + rmm::device_uvector merged_majors(0, handle.get_stream()); + rmm::device_uvector merged_minors(0, handle.get_stream()); + std::vector merged_properties(lower_triangular_properties.size(), + handle.get_stream()); + + std::vector properties_in_unique_key{}; + for (size_t i = 0; i < edge_property_combining_rules.size(); ++i) { + if (edge_property_combining_rules[i] == edge_property_combining_rule_t::PART_OF_UNIQUENESS) { + properties_in_unique_key.push_back(i); + } + } + + if (properties_in_unique_key.size() == 0) { + thrust::sort( + handle.get_thrust_policy(), + thrust::make_zip_iterator(lower_triangular_majors.begin(), lower_triangular_minors.begin()), + thrust::make_zip_iterator(lower_triangular_majors.end(), lower_triangular_minors.end())); + thrust::sort(handle.get_thrust_policy(), + upper_triangular_edge_first, + upper_triangular_edge_first + upper_triangular_majors.size()); + } else if (properties_in_unique_key.size() == 1) { + cugraph::variant_type_dispatch( + lower_triangular_properties[properties_in_unique_key[0]], + [&lower_triangular_majors, &lower_triangular_minors, &lower_triangular_properties]( + auto& lower_triangular_property) { + using T = std::decay_t; + + thrust::sort(handle.get_thrust_policy(), + lower_triangular_edge_first, + lower_triangular_edge_first + lower_triangular_majors.size(), + compare_lower_triangular_edges_as_lower_triangular_t{}); + + thrust::sort(handle.get_thrust_policy(), + upper_triangular_edge_first, + upper_triangular_edge_first + upper_triangular_majors.size(), + compare_upper_triangular_edges_as_lower_triangular_t{}); + }); + thrust::sort(handle.get_thrust_policy(), + lower_triangular_edge_first, + lower_triangular_edge_first + lower_triangular_majors.size(), + compare_lower_triangular_edges_as_lower_triangular_t{}); + thrust::sort(handle.get_thrust_policy(), + upper_triangular_edge_first, + upper_triangular_edge_first + upper_triangular_majors.size(), + compare_upper_triangular_edges_as_lower_triangular_t{}); + } else { + CUGRAPH_FAIL("PART_OF_UNIQUENESS is supported for at most one property"); + } + + auto lower_triangular_edge_first = thrust::make_zip_iterator(lower_triangular_majors.begin(), + lower_triangular_minors.begin(), + lower_triangular_properties.begin()); + thrust::sort(handle.get_thrust_policy(), + lower_triangular_edge_first, + lower_triangular_edge_first + lower_triangular_majors.size()); + auto upper_triangular_edge_first = thrust::make_zip_iterator( + upper_triangular_majors.begin(), + upper_triangular_minors.begin(), + upper_triangular_properties + .begin()); // do not flip here to use "lower_triangular = major > minor" + thrust::sort(handle.get_thrust_policy(), + upper_triangular_edge_first, + upper_triangular_edge_first + upper_triangular_majors.size(), + compare_upper_triangular_edges_as_lower_triangular_t{}); + + merged_majors.resize(lower_triangular_majors.size() + upper_triangular_majors.size(), + handle.get_stream()); + merged_minors.resize(merged_majors.size(), handle.get_stream()); + merged_properties.resize(merged_majors.size(), handle.get_stream()); + auto merged_first = thrust::make_zip_iterator( + merged_majors.begin(), merged_minors.begin(), merged_properties.begin()); + thrust::merge(handle.get_thrust_policy(), + lower_triangular_edge_first, + lower_triangular_edge_first + lower_triangular_majors.size(), + upper_triangular_edge_first, + upper_triangular_edge_first + upper_triangular_majors.size(), + merged_first, + compare_lower_and_upper_triangular_edges_t{}); + + lower_triangular_majors.resize(0, handle.get_stream()); + lower_triangular_majors.shrink_to_fit(handle.get_stream()); + lower_triangular_minors.resize(0, handle.get_stream()); + lower_triangular_minors.shrink_to_fit(handle.get_stream()); + lower_triangular_properties.resize(0, handle.get_stream()); + lower_triangular_properties.shrink_to_fit(handle.get_stream()); + + upper_triangular_majors.resize(0, handle.get_stream()); + upper_triangular_majors.shrink_to_fit(handle.get_stream()); + upper_triangular_minors.resize(0, handle.get_stream()); + upper_triangular_minors.shrink_to_fit(handle.get_stream()); + upper_triangular_properties.resize(0, handle.get_stream()); + upper_triangular_properties.shrink_to_fit(handle.get_stream()); + + rmm::device_uvector includes(merged_majors.size(), handle.get_stream()); + symmetrize_op_t symm_op{reciprocal}; + thrust::for_each(handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(merged_majors.size()), + update_edge_weights_and_flags_t{ + merged_first, includes.data(), merged_majors.size(), symm_op}); + + auto merged_edge_and_flag_first = thrust::make_zip_iterator( + merged_majors.begin(), merged_minors.begin(), merged_properties.begin(), includes.begin()); + merged_majors.resize( + cuda::std::distance(merged_edge_and_flag_first, + thrust::remove_if(handle.get_thrust_policy(), + merged_edge_and_flag_first, + merged_edge_and_flag_first + merged_majors.size(), + [] __device__(auto t) { return !cuda::std::get<3>(t); })), + handle.get_stream()); + merged_majors.shrink_to_fit(handle.get_stream()); + merged_minors.resize(merged_majors.size(), handle.get_stream()); + merged_minors.shrink_to_fit(handle.get_stream()); + merged_properties.resize(merged_majors.size(), handle.get_stream()); + merged_properties.shrink_to_fit(handle.get_stream()); + + auto merged_major_minor_range_first = + thrust::make_zip_iterator(merged_majors.begin(), merged_minors.begin()); + thrust::transform(handle.get_thrust_policy(), + merged_major_minor_range_first, + merged_major_minor_range_first + merged_majors.size(), + merged_major_minor_range_first, + to_lower_triangular_t{}); + + return std::make_tuple( + std::move(merged_majors), std::move(merged_minors), std::move(merged_properties)); +} + template std::tuple, rmm::device_uvector, @@ -346,35 +490,19 @@ std::tuple, rmm::device_uvector> merge_l return std::make_tuple(std::move(merged_majors), std::move(merged_minors)); } -template +template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -symmetrize_edgelist(raft::handle_t const& handle, - rmm::device_uvector&& edgelist_majors, - rmm::device_uvector&& edgelist_minors, - std::optional>&& edgelist_weights, - std::optional>&& edgelist_edge_ids, - std::optional>&& edgelist_edge_types, - std::optional>&& edgelist_edge_start_times, - std::optional>&& edgelist_edge_end_times, - bool reciprocal) + std::vector> +symmetrize_edgelist( + raft::handle_t const& handle, + rmm::device_uvector&& edgelist_majors, + rmm::device_uvector&& edgelist_minors, + std::vector&& edgelist_edge_properties, + std::vector const& edge_property_combining_rules, + bool remove_edges) { - int edge_property_count = 0; - if (edgelist_weights) ++edge_property_count; - if (edgelist_edge_ids) ++edge_property_count; - if (edgelist_edge_types) ++edge_property_count; - if (edgelist_edge_start_times) ++edge_property_count; - if (edgelist_edge_end_times) ++edge_property_count; + int edge_property_count = edgelist_edge_properties.size(); // 1. separate lower triangular, diagonal (self-loop), and upper triangular edges @@ -408,90 +536,28 @@ symmetrize_edgelist(raft::handle_t const& handle, num_diagonal_edges = static_cast(cuda::std::distance(lower_triangular_last, diagonal_last)); } else if (edge_property_count == 1) { - if (edgelist_weights) { - auto edge_first = thrust::make_zip_iterator( - edgelist_majors.begin(), edgelist_minors.begin(), edgelist_weights->begin()); - auto lower_triangular_last = thrust::partition(handle.get_thrust_policy(), - edge_first, - edge_first + edgelist_majors.size(), - lower_triangular_compare); - num_lower_triangular_edges = - static_cast(cuda::std::distance(edge_first, lower_triangular_last)); - auto diagonal_last = thrust::partition(handle.get_thrust_policy(), - edge_first + num_lower_triangular_edges, - edge_first + edgelist_majors.size(), - diagonal_compare); - num_diagonal_edges = - static_cast(cuda::std::distance(lower_triangular_last, diagonal_last)); - } - - if (edgelist_edge_ids) { - auto edge_first = thrust::make_zip_iterator( - edgelist_majors.begin(), edgelist_minors.begin(), edgelist_edge_ids->begin()); - auto lower_triangular_last = thrust::partition(handle.get_thrust_policy(), - edge_first, - edge_first + edgelist_majors.size(), - lower_triangular_compare); - num_lower_triangular_edges = - static_cast(cuda::std::distance(edge_first, lower_triangular_last)); - auto diagonal_last = thrust::partition(handle.get_thrust_policy(), - edge_first + num_lower_triangular_edges, - edge_first + edgelist_majors.size(), - diagonal_compare); - num_diagonal_edges = - static_cast(cuda::std::distance(lower_triangular_last, diagonal_last)); - } - - if (edgelist_edge_types) { - auto edge_first = thrust::make_zip_iterator( - edgelist_majors.begin(), edgelist_minors.begin(), edgelist_edge_types->begin()); - auto lower_triangular_last = thrust::partition(handle.get_thrust_policy(), - edge_first, - edge_first + edgelist_majors.size(), - lower_triangular_compare); - num_lower_triangular_edges = - static_cast(cuda::std::distance(edge_first, lower_triangular_last)); - auto diagonal_last = thrust::partition(handle.get_thrust_policy(), - edge_first + num_lower_triangular_edges, - edge_first + edgelist_majors.size(), - diagonal_compare); - num_diagonal_edges = - static_cast(cuda::std::distance(lower_triangular_last, diagonal_last)); - } - - if (edgelist_edge_start_times) { - auto edge_first = thrust::make_zip_iterator( - edgelist_majors.begin(), edgelist_minors.begin(), edgelist_edge_start_times->begin()); - auto lower_triangular_last = thrust::partition(handle.get_thrust_policy(), - edge_first, - edge_first + edgelist_majors.size(), - lower_triangular_compare); - num_lower_triangular_edges = - static_cast(cuda::std::distance(edge_first, lower_triangular_last)); - auto diagonal_last = thrust::partition(handle.get_thrust_policy(), - edge_first + num_lower_triangular_edges, - edge_first + edgelist_majors.size(), - diagonal_compare); - num_diagonal_edges = - static_cast(cuda::std::distance(lower_triangular_last, diagonal_last)); - } - - if (edgelist_edge_end_times) { - auto edge_first = thrust::make_zip_iterator( - edgelist_majors.begin(), edgelist_minors.begin(), edgelist_edge_end_times->begin()); - auto lower_triangular_last = thrust::partition(handle.get_thrust_policy(), - edge_first, - edge_first + edgelist_majors.size(), - lower_triangular_compare); - num_lower_triangular_edges = - static_cast(cuda::std::distance(edge_first, lower_triangular_last)); - auto diagonal_last = thrust::partition(handle.get_thrust_policy(), - edge_first + num_lower_triangular_edges, - edge_first + edgelist_majors.size(), - diagonal_compare); - num_diagonal_edges = - static_cast(cuda::std::distance(lower_triangular_last, diagonal_last)); - } + std::tie(num_lower_triangular_edges, num_diagonal_edges) = cugraph::variant_type_dispatch( + edgelist_edge_properties[0], + [&edgelist_majors, &edgelist_minors, &lower_triangular_compare, &diagonal_compare]( + auto& edgelist_edge_property) { + using T = std::decay_t; + + auto edge_first = thrust::make_zip_iterator( + edgelist_majors.begin(), edgelist_minors.begin(), edgelist_edge_property.begin()); + auto lower_triangular_last = thrust::partition(handle.get_thrust_policy(), + edge_first, + edge_first + edgelist_majors.size(), + lower_triangular_compare); + num_lower_triangular_edges = + static_cast(cuda::std::distance(edge_first, lower_triangular_last)); + auto diagonal_last = thrust::partition(handle.get_thrust_policy(), + edge_first + num_lower_triangular_edges, + edge_first + edgelist_majors.size(), + diagonal_compare); + num_diagonal_edges = + static_cast(cuda::std::distance(lower_triangular_last, diagonal_last)); + return std::make_tuple(num_lower_triangular_edges, num_diagonal_edges); + }); } else { rmm::device_uvector property_position(edgelist_majors.size(), handle.get_stream()); detail::sequence_fill( @@ -512,64 +578,33 @@ symmetrize_edgelist(raft::handle_t const& handle, num_diagonal_edges = static_cast(cuda::std::distance(lower_triangular_last, diagonal_last)); - if (edgelist_weights) { - rmm::device_uvector tmp(property_position.size(), handle.get_stream()); - - thrust::gather(handle.get_thrust_policy(), - property_position.begin(), - property_position.end(), - edgelist_weights->begin(), - tmp.begin()); - edgelist_weights = std::move(tmp); - } - - if (edgelist_edge_ids) { - rmm::device_uvector tmp(property_position.size(), handle.get_stream()); - - thrust::gather(handle.get_thrust_policy(), - property_position.begin(), - property_position.end(), - edgelist_edge_ids->begin(), - tmp.begin()); - - edgelist_edge_ids = std::move(tmp); - } - - if (edgelist_edge_types) { - rmm::device_uvector tmp(property_position.size(), handle.get_stream()); - - thrust::gather(handle.get_thrust_policy(), - property_position.begin(), - property_position.end(), - edgelist_edge_types->begin(), - tmp.begin()); - - edgelist_edge_types = std::move(tmp); - } - - if (edgelist_edge_start_times) { - rmm::device_uvector tmp(property_position.size(), handle.get_stream()); - - thrust::gather(handle.get_thrust_policy(), - property_position.begin(), - property_position.end(), - edgelist_edge_start_times->begin(), - tmp.begin()); - - edgelist_edge_start_times = std::move(tmp); - } - - if (edgelist_edge_end_times) { - rmm::device_uvector tmp(property_position.size(), handle.get_stream()); - - thrust::gather(handle.get_thrust_policy(), - property_position.begin(), - property_position.end(), - edgelist_edge_end_times->begin(), - tmp.begin()); - - edgelist_edge_end_times = std::move(tmp); - } + std::for_each(edgelist_edge_properties.begin(), + edgelist_edge_properties.end(), + [&property_position, + &edgelist_majors, + &edgelist_minors, + &lower_triangular_compare, + &diagonal_compare](auto& edgelist_edge_property) { + cugraph::variant_type_dispatch( + edgelist_edge_property, + [&property_position, + &edgelist_majors, + &edgelist_minors, + &lower_triangular_compare, + &diagonal_compare](auto& edgelist_edge_property) { + using T = std::decay_t; + + rmm::device_uvector tmp(property_position.size(), handle.get_stream()); + + thrust::gather(handle.get_thrust_policy(), + property_position.begin(), + property_position.end(), + edgelist_edge_property.begin(), + tmp.begin()); + + edgelist_edge_property = std::move(tmp); + }); + }); } rmm::device_uvector diagonal_majors(num_diagonal_edges, handle.get_stream()); @@ -597,121 +632,55 @@ symmetrize_edgelist(raft::handle_t const& handle, edgelist_minors.shrink_to_fit(handle.get_stream()); auto lower_triangular_minors = std::move(edgelist_minors); - std::optional> diagonal_weights{std::nullopt}; - std::optional> diagonal_edge_ids{std::nullopt}; - std::optional> diagonal_edge_types{std::nullopt}; - std::optional> diagonal_edge_start_times{std::nullopt}; - std::optional> diagonal_edge_end_times{std::nullopt}; - std::optional> upper_triangular_weights{std::nullopt}; - std::optional> upper_triangular_edge_ids{std::nullopt}; - std::optional> upper_triangular_edge_types{std::nullopt}; - std::optional> upper_triangular_edge_start_times{std::nullopt}; - std::optional> upper_triangular_edge_end_times{std::nullopt}; - - if (edgelist_weights) { - diagonal_weights = rmm::device_uvector(num_diagonal_edges, handle.get_stream()); - upper_triangular_weights = - rmm::device_uvector(upper_triangular_majors.size(), handle.get_stream()); - thrust::copy(handle.get_thrust_policy(), - edgelist_weights->begin() + num_lower_triangular_edges, - edgelist_weights->begin() + num_lower_triangular_edges + num_diagonal_edges, - diagonal_weights->begin()); - thrust::copy(handle.get_thrust_policy(), - edgelist_weights->begin() + num_lower_triangular_edges + num_diagonal_edges, - edgelist_weights->end(), - upper_triangular_weights->begin()); - edgelist_weights->resize(lower_triangular_majors.size(), handle.get_stream()); - edgelist_weights->shrink_to_fit(handle.get_stream()); - } - auto lower_triangular_weights = std::move(edgelist_weights); - - if (edgelist_edge_ids) { - diagonal_edge_ids = rmm::device_uvector(num_diagonal_edges, handle.get_stream()); - upper_triangular_edge_ids = - rmm::device_uvector(upper_triangular_majors.size(), handle.get_stream()); - thrust::copy(handle.get_thrust_policy(), - edgelist_edge_ids->begin() + num_lower_triangular_edges, - edgelist_edge_ids->begin() + num_lower_triangular_edges + num_diagonal_edges, - diagonal_edge_ids->begin()); - thrust::copy(handle.get_thrust_policy(), - edgelist_edge_ids->begin() + num_lower_triangular_edges + num_diagonal_edges, - edgelist_edge_ids->end(), - upper_triangular_edge_ids->begin()); - edgelist_edge_ids->resize(lower_triangular_majors.size(), handle.get_stream()); - edgelist_edge_ids->shrink_to_fit(handle.get_stream()); - } - auto lower_triangular_edge_ids = std::move(edgelist_edge_ids); - - if (edgelist_edge_types) { - diagonal_edge_types = rmm::device_uvector(num_diagonal_edges, handle.get_stream()); - upper_triangular_edge_types = - rmm::device_uvector(upper_triangular_majors.size(), handle.get_stream()); - thrust::copy(handle.get_thrust_policy(), - edgelist_edge_types->begin() + num_lower_triangular_edges, - edgelist_edge_types->begin() + num_lower_triangular_edges + num_diagonal_edges, - diagonal_edge_types->begin()); - thrust::copy(handle.get_thrust_policy(), - edgelist_edge_types->begin() + num_lower_triangular_edges + num_diagonal_edges, - edgelist_edge_types->end(), - upper_triangular_edge_types->begin()); - edgelist_edge_types->resize(lower_triangular_majors.size(), handle.get_stream()); - edgelist_edge_types->shrink_to_fit(handle.get_stream()); - } - auto lower_triangular_edge_types = std::move(edgelist_edge_types); - - if (edgelist_edge_start_times) { - diagonal_edge_start_times = - rmm::device_uvector(num_diagonal_edges, handle.get_stream()); - upper_triangular_edge_start_times = - rmm::device_uvector(upper_triangular_majors.size(), handle.get_stream()); - thrust::copy( - handle.get_thrust_policy(), - edgelist_edge_start_times->begin() + num_lower_triangular_edges, - edgelist_edge_start_times->begin() + num_lower_triangular_edges + num_diagonal_edges, - diagonal_edge_start_times->begin()); - thrust::copy( - handle.get_thrust_policy(), - edgelist_edge_start_times->begin() + num_lower_triangular_edges + num_diagonal_edges, - edgelist_edge_start_times->end(), - upper_triangular_edge_start_times->begin()); - edgelist_edge_start_times->resize(lower_triangular_majors.size(), handle.get_stream()); - edgelist_edge_start_times->shrink_to_fit(handle.get_stream()); - } - auto lower_triangular_edge_start_times = std::move(edgelist_edge_start_times); - - if (edgelist_edge_end_times) { - diagonal_edge_end_times = - rmm::device_uvector(num_diagonal_edges, handle.get_stream()); - upper_triangular_edge_end_times = - rmm::device_uvector(upper_triangular_majors.size(), handle.get_stream()); - thrust::copy(handle.get_thrust_policy(), - edgelist_edge_end_times->begin() + num_lower_triangular_edges, - edgelist_edge_end_times->begin() + num_lower_triangular_edges + num_diagonal_edges, - diagonal_edge_end_times->begin()); - thrust::copy(handle.get_thrust_policy(), - edgelist_edge_end_times->begin() + num_lower_triangular_edges + num_diagonal_edges, - edgelist_edge_end_times->end(), - upper_triangular_edge_end_times->begin()); - edgelist_edge_end_times->resize(lower_triangular_majors.size(), handle.get_stream()); - edgelist_edge_end_times->shrink_to_fit(handle.get_stream()); - } - auto lower_triangular_edge_end_times = std::move(edgelist_edge_end_times); + std::vector lower_triangular_edge_properties(edge_property_count); + std::vector diagonal_edge_properties(edge_property_count); + std::vector upper_triangular_edge_properties(edge_property_count); + + std::for_each( + edgelist_edge_properties.begin(), + edgelist_edge_properties.end(), + [&lower_triangular_edge_properties, + &diagonal_edge_properties, + &upper_triangular_edge_properties, + num_diagonal_edges, + num_upper_triangular_edges = upper_triangular_majors.size(), + num_lower_triangular_edges](auto& edgelist_edge_property) { + cugraph::variant_type_dispatch( + edgelist_edge_property, + [&lower_triangular_edge_properties, + &diagonal_edge_properties, + &upper_triangular_edge_properties, + num_diagonal_edges, + num_upper_triangular_edges, + num_lower_triangular_edges](auto& edgelist_edge_property) { + using T = std::decay_t; + + rmm::device_uvector diagonal_tmp(num_diagonal_edges, handle.get_stream()); + rmm::device_uvector upper_triangular_tmp(num_upper_triangular_edges, + handle.get_stream()); + + thrust::copy( + handle.get_thrust_policy(), + edgelist_edge_property.begin() + num_lower_triangular_edges, + edgelist_edge_property.begin() + num_lower_triangular_edges + num_diagonal_edges, + diagonal_tmp.begin()); + thrust::copy( + handle.get_thrust_policy(), + edgelist_edge_property.begin() + num_lower_triangular_edges + num_diagonal_edges, + edgelist_edge_property.end(), + upper_triangular_tmp.begin()); + edgelist_edge_property.resize(lower_triangular_majors.size(), handle.get_stream()); + edgelist_edge_property.shrink_to_fit(handle.get_stream()); + + lower_triangular_edge_properties.push_back(std::move(edgelist_edge_property)); + diagonal_edge_properties.push_back(std::move(diagonal_tmp)); + upper_triangular_edge_properties.push_back(std::move(upper_triangular_tmp)); + }); + }); // 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_edge_properties, @@ -720,47 +689,14 @@ symmetrize_edgelist(raft::handle_t const& handle, 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 rmm::device_uvector merged_lower_triangular_majors(0, handle.get_stream()); rmm::device_uvector merged_lower_triangular_minors(0, handle.get_stream()); - auto merged_lower_triangular_weights = - edgelist_weights ? std::make_optional>(0, handle.get_stream()) - : std::nullopt; - auto merged_lower_triangular_edge_ids = - edgelist_edge_ids ? std::make_optional>(0, handle.get_stream()) - : std::nullopt; - auto merged_lower_triangular_edge_types = - edgelist_edge_types - ? std::make_optional>(0, handle.get_stream()) - : std::nullopt; - auto merged_lower_triangular_edge_start_times = - edgelist_edge_start_times - ? std::make_optional>(0, handle.get_stream()) - : std::nullopt; - auto merged_lower_triangular_edge_end_times = - edgelist_edge_end_times - ? std::make_optional>(0, handle.get_stream()) - : std::nullopt; + std::vector merged_lower_triangular_edge_properties( + edge_property_count); if (edge_property_count == 0) { std::tie(merged_lower_triangular_majors, merged_lower_triangular_minors) = @@ -770,21 +706,37 @@ symmetrize_edgelist(raft::handle_t const& handle, std::move(upper_triangular_majors), std::move(upper_triangular_minors), num_lower_triangular_edges, - reciprocal); + remove_edges); } else if (edge_property_count == 1) { + cugraph::variant_type_dispatch( + lower_triangular_edge_properties[0], + [&merged_lower_triangular_edge_properties, + &merged_lower_triangular_majors, + &merged_lower_triangular_minors, + &edge_property_combining_rules, + &upper_triangular_edge_properties](auto& edgelist_edge_property) { + using T = std::decay_t; + + rmm::device_uvector merged_lower_triangular_tmp(merged_lower_triangular_majors.size(), + handle.get_stream()); + + std::tie(merged_lower_triangular_majors, + merged_lower_triangular_minors, + merged_lower_triangular_edge_property) = + merge_lower_triangular( + handle, + std::move(lower_triangular_majors), + std::move(lower_triangular_minors), + std::move(edgelist_edge_property), + std::move(upper_triangular_majors), + std::move(upper_triangular_minors), + std::move(std::get>(upper_triangular_edge_properties[0])), + edge_property_combining_rules, + num_lower_triangular_edges, + remove_edges); + }); + if (edgelist_weights) { - std::tie(merged_lower_triangular_majors, - merged_lower_triangular_minors, - merged_lower_triangular_weights) = - merge_lower_triangular(handle, - std::move(lower_triangular_majors), - std::move(lower_triangular_minors), - std::move(*lower_triangular_weights), - std::move(upper_triangular_majors), - std::move(upper_triangular_minors), - std::move(*upper_triangular_weights), - num_lower_triangular_edges, - reciprocal); } else if (edgelist_edge_ids) { std::tie(merged_lower_triangular_majors, merged_lower_triangular_minors, @@ -797,7 +749,7 @@ symmetrize_edgelist(raft::handle_t const& handle, std::move(upper_triangular_minors), std::move(*upper_triangular_edge_ids), num_lower_triangular_edges, - reciprocal); + remove_edges); } else if (edgelist_edge_types) { std::tie(merged_lower_triangular_majors, merged_lower_triangular_minors, @@ -810,7 +762,7 @@ symmetrize_edgelist(raft::handle_t const& handle, std::move(upper_triangular_minors), std::move(*upper_triangular_edge_types), num_lower_triangular_edges, - reciprocal); + remove_edges); } else if (edgelist_edge_start_times) { std::tie(merged_lower_triangular_majors, merged_lower_triangular_minors, @@ -823,7 +775,7 @@ symmetrize_edgelist(raft::handle_t const& handle, std::move(upper_triangular_minors), std::move(*upper_triangular_edge_start_times), num_lower_triangular_edges, - reciprocal); + remove_edges); } else if (edgelist_edge_end_times) { std::tie(merged_lower_triangular_majors, merged_lower_triangular_minors, @@ -836,7 +788,7 @@ symmetrize_edgelist(raft::handle_t const& handle, std::move(upper_triangular_minors), std::move(*upper_triangular_edge_end_times), num_lower_triangular_edges, - reciprocal); + remove_edges); } } else { rmm::device_uvector property_position( @@ -853,7 +805,7 @@ symmetrize_edgelist(raft::handle_t const& handle, std::move(upper_triangular_minors), std::move(property_position), num_lower_triangular_edges, - reciprocal); + remove_edges); if (edgelist_weights) { merged_lower_triangular_weights->resize(merged_lower_triangular_majors.size(), @@ -1162,59 +1114,64 @@ symmetrize_edgelist(raft::handle_t const& handle, } // namespace detail -template +template std::tuple, rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -symmetrize_edgelist(raft::handle_t const& handle, - rmm::device_uvector&& edgelist_srcs, - rmm::device_uvector&& edgelist_dsts, - std::optional>&& edgelist_weights, - std::optional>&& edgelist_edge_ids, - std::optional>&& edgelist_edge_types, - std::optional>&& edgelist_edge_start_times, - std::optional>&& edgelist_edge_end_times, - bool reciprocal) + std::vector> +add_reciprocal_edges( + raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::vector&& edgelist_edge_properties, + std::vector const& edge_property_combining_rules, + bool store_transposed) { auto edgelist_majors = std::move(store_transposed ? edgelist_dsts : edgelist_srcs); auto edgelist_minors = std::move(store_transposed ? edgelist_srcs : edgelist_dsts); - std::tie(edgelist_majors, - edgelist_minors, - edgelist_weights, - edgelist_edge_ids, - edgelist_edge_types, - edgelist_edge_start_times, - edgelist_edge_end_times) = - detail::symmetrize_edgelist( + + std::tie(edgelist_majors, edgelist_minors, edgelist_edge_properties) = + detail::symmetrize_edgelist( handle, std::move(edgelist_majors), std::move(edgelist_minors), - 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), - reciprocal); + std::move(edgelist_edge_properties), + std::vector{}, //? + false); + edgelist_srcs = std::move(store_transposed ? edgelist_minors : edgelist_majors); edgelist_dsts = std::move(store_transposed ? edgelist_majors : edgelist_minors); - return std::make_tuple(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)); + return std::make_tuple( + std::move(edgelist_srcs), std::move(edgelist_dsts), std::move(edgelist_edge_properties)); +} + +template +std::tuple, + rmm::device_uvector, + std::vector> +remove_nonreciprocal_edges(raft::handle_t const& handle, + rmm::device_uvector&& edgelist_srcs, + rmm::device_uvector&& edgelist_dsts, + std::vector&& edgelist_edge_properties, + bool store_transposed) +{ + auto edgelist_majors = std::move(store_transposed ? edgelist_dsts : edgelist_srcs); + auto edgelist_minors = std::move(store_transposed ? edgelist_srcs : edgelist_dsts); + + std::tie(edgelist_majors, edgelist_minors, edgelist_edge_properties) = + detail::symmetrize_edgelist( + handle, + std::move(edgelist_majors), + std::move(edgelist_minors), + std::move(edgelist_edge_properties), + std::vector{}, //? + true); + + edgelist_srcs = std::move(store_transposed ? edgelist_minors : edgelist_majors); + edgelist_dsts = std::move(store_transposed ? edgelist_majors : edgelist_minors); + + return std::make_tuple( + std::move(edgelist_srcs), std::move(edgelist_dsts), std::move(edgelist_edge_properties)); } } // namespace cugraph