diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index d3dfdbd068c..57e0aa2d078 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -288,6 +288,8 @@ set(CUGRAPH_SOURCES src/structure/symmetrize_edgelist_mg.cu src/community/triangle_count_sg.cu src/community/triangle_count_mg.cu + src/community/approx_weighted_matching_sg.cu + src/community/approx_weighted_matching_mg.cu src/traversal/k_hop_nbrs_sg.cu src/traversal/k_hop_nbrs_mg.cu src/mtmg/vertex_result.cu diff --git a/cpp/include/cugraph/algorithms.hpp b/cpp/include/cugraph/algorithms.hpp index 0caa151daac..7c4a978c4b4 100644 --- a/cpp/include/cugraph/algorithms.hpp +++ b/cpp/include/cugraph/algorithms.hpp @@ -2368,6 +2368,32 @@ rmm::device_uvector vertex_coloring( graph_view_t const& graph_view, raft::random::RngState& rng_state); +/* + * @brief Approximate Weighted Matching + * + * A matching in an undirected graph G = (V, E) is a pairing of adjacent vertices + * such that each vertex is matched with at most one other vertex, the objective + * being to match as many vertices as possible or to maximise the sum of the + * weights of the matched edges. Here we provide an implementation of an + * approximation algorithm to the weighted Maximum matching. See + * https://web.archive.org/web/20081031230449id_/http://www.ii.uib.no/~fredrikm/fredrik/papers/CP75.pdf + * for further information. + * + * @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 multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * @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] graph_view Graph view object. + * @param[in] edge_weight_view View object holding edge weights for @p graph_view. + * @return A tuple of device vector of matched vertex ids and sum of the weights of the matched + * edges. + */ +template +std::tuple, weight_t> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); } // namespace cugraph /** diff --git a/cpp/src/community/approx_weighted_matching_impl.cuh b/cpp/src/community/approx_weighted_matching_impl.cuh new file mode 100644 index 00000000000..e693beee489 --- /dev/null +++ b/cpp/src/community/approx_weighted_matching_impl.cuh @@ -0,0 +1,392 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#pragma once + +#include "prims/fill_edge_property.cuh" +#include "prims/reduce_op.cuh" +#include "prims/transform_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 + +#include + +namespace cugraph { + +namespace detail { + +template +std::tuple, weight_t> approximate_weighted_matching( + raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + edge_property_view_t edge_weight_view) +{ + CUGRAPH_EXPECTS(graph_view.is_symmetric(), + "Invalid input arguments: input graph for approximate_weighted_matching must " + "need to be symmetric"); + + using graph_view_t = cugraph::graph_view_t; + + graph_view_t current_graph_view(graph_view); + if (current_graph_view.has_edge_mask()) { current_graph_view.clear_edge_mask(); } + + cugraph::edge_property_t edge_masks_even(handle, current_graph_view); + cugraph::fill_edge_property(handle, current_graph_view, bool{false}, edge_masks_even); + cugraph::edge_property_t edge_masks_odd(handle, current_graph_view); + cugraph::fill_edge_property(handle, current_graph_view, bool{false}, edge_masks_odd); + + if (graph_view.has_edge_mask()) { + current_graph_view.attach_edge_mask(*(graph_view.edge_mask_view())); + } + // Mask out self-loop + cugraph::transform_e( + handle, + current_graph_view, + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + cugraph::edge_dummy_property_t{}.view(), + [] __device__(auto src, auto dst, thrust::nullopt_t, thrust::nullopt_t, thrust::nullopt_t) { + return !(src == dst); + }, + edge_masks_even.mutable_view()); + + if (current_graph_view.has_edge_mask()) current_graph_view.clear_edge_mask(); + current_graph_view.attach_edge_mask(edge_masks_even.view()); + + auto constexpr invalid_partner = invalid_vertex_id::value; + rmm::device_uvector offers_from_partners( + current_graph_view.local_vertex_partition_range_size(), handle.get_stream()); + + rmm::device_uvector partners(current_graph_view.local_vertex_partition_range_size(), + handle.get_stream()); + + thrust::fill(handle.get_thrust_policy(), partners.begin(), partners.end(), invalid_partner); + thrust::fill(handle.get_thrust_policy(), + offers_from_partners.begin(), + offers_from_partners.end(), + weight_t{0.0}); + + rmm::device_uvector local_vertices( + current_graph_view.local_vertex_partition_range_size(), handle.get_stream()); + detail::sequence_fill(handle.get_stream(), + local_vertices.begin(), + local_vertices.size(), + current_graph_view.local_vertex_partition_range_first()); + + edge_src_property_t src_key_cache(handle); + cugraph::edge_src_property_t src_match_flags(handle); + cugraph::edge_dst_property_t dst_match_flags(handle); + + if constexpr (graph_view_t::is_multi_gpu) { + src_key_cache = edge_src_property_t(handle, current_graph_view); + + update_edge_src_property(handle, current_graph_view, local_vertices.begin(), src_key_cache); + + src_match_flags = cugraph::edge_src_property_t(handle, current_graph_view); + dst_match_flags = cugraph::edge_dst_property_t(handle, current_graph_view); + } + + vertex_t loop_counter = 0; + while (true) { + // + // For each candidate vertex, find the best possible target + // + + rmm::device_uvector candidates(0, handle.get_stream()); + rmm::device_uvector offers_from_candidates(0, handle.get_stream()); + rmm::device_uvector targets(0, handle.get_stream()); + + // FIXME: This can be implemented more efficiently if per_v_transform_reduce_incoming|outgoing_e + // is updated to support reduction on thrust::tuple. + std::forward_as_tuple(candidates, std::tie(offers_from_candidates, targets)) = + cugraph::transform_reduce_e_by_src_key( + handle, + current_graph_view, + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + edge_weight_view, + graph_view_t::is_multi_gpu + ? src_key_cache.view() + : detail::edge_major_property_view_t(local_vertices.begin()), + [] __device__(auto, auto dst, thrust::nullopt_t, thrust::nullopt_t, auto wt) { + return thrust::make_tuple(wt, dst); + }, + thrust::make_tuple(weight_t{0.0}, invalid_partner), + reduce_op::maximum>{}, + true); + + // + // For each target, find the best offer + // + + if constexpr (graph_view_t::is_multi_gpu) { + auto vertex_partition_range_lasts = current_graph_view.vertex_partition_range_lasts(); + + rmm::device_uvector d_vertex_partition_range_lasts( + vertex_partition_range_lasts.size(), handle.get_stream()); + + raft::update_device(d_vertex_partition_range_lasts.data(), + vertex_partition_range_lasts.data(), + vertex_partition_range_lasts.size(), + handle.get_stream()); + + 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(); + + auto key_func = cugraph::detail::compute_gpu_id_from_int_vertex_t{ + raft::device_span(d_vertex_partition_range_lasts.data(), + d_vertex_partition_range_lasts.size()), + major_comm_size, + minor_comm_size}; + + std::forward_as_tuple(std::tie(candidates, offers_from_candidates, targets), std::ignore) = + cugraph::groupby_gpu_id_and_shuffle_values( + handle.get_comms(), + thrust::make_zip_iterator(thrust::make_tuple( + candidates.begin(), offers_from_candidates.begin(), targets.begin())), + thrust::make_zip_iterator( + thrust::make_tuple(candidates.end(), offers_from_candidates.end(), targets.end())), + [key_func] __device__(auto val) { return key_func(thrust::get<2>(val)); }, + handle.get_stream()); + } + + auto itr_to_tuples = thrust::make_zip_iterator( + thrust::make_tuple(offers_from_candidates.begin(), candidates.begin())); + + thrust::sort_by_key(handle.get_thrust_policy(), targets.begin(), targets.end(), itr_to_tuples); + + auto nr_unique_targets = thrust::count_if(handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(targets.size()), + is_first_in_run_t{targets.data()}); + + rmm::device_uvector unique_targets(nr_unique_targets, handle.get_stream()); + rmm::device_uvector best_offers_to_targets(nr_unique_targets, handle.get_stream()); + rmm::device_uvector best_candidates(nr_unique_targets, handle.get_stream()); + + auto itr_to_reduced_tuples = thrust::make_zip_iterator( + thrust::make_tuple(best_offers_to_targets.begin(), best_candidates.begin())); + + auto new_end = thrust::reduce_by_key( + handle.get_thrust_policy(), + targets.begin(), + targets.end(), + itr_to_tuples, + unique_targets.begin(), + itr_to_reduced_tuples, + thrust::equal_to{}, + [] __device__(auto pair1, auto pair2) { return (pair1 > pair2) ? pair1 : pair2; }); + + vertex_t nr_reduces_tuples = + static_cast(thrust::distance(unique_targets.begin(), new_end.first)); + + targets = std::move(unique_targets); + offers_from_candidates = std::move(best_offers_to_targets); + candidates = std::move(best_candidates); + + // + // two vertex offer each other, that's a match + // + + kv_store_t target_candidate_map(targets.begin(), + targets.end(), + candidates.begin(), + invalid_vertex_id::value, + invalid_vertex_id::value, + handle.get_stream()); + + rmm::device_uvector candidates_of_candidates(0, handle.get_stream()); + + if (graph_view_t::is_multi_gpu) { + auto& comm = handle.get_comms(); + 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(); + + auto partitions_range_lasts = graph_view.vertex_partition_range_lasts(); + rmm::device_uvector d_partitions_range_lasts(partitions_range_lasts.size(), + handle.get_stream()); + + raft::update_device(d_partitions_range_lasts.data(), + partitions_range_lasts.data(), + partitions_range_lasts.size(), + handle.get_stream()); + + cugraph::detail::compute_gpu_id_from_int_vertex_t vertex_to_gpu_id_op{ + raft::device_span(d_partitions_range_lasts.data(), + d_partitions_range_lasts.size()), + major_comm_size, + minor_comm_size}; + + candidates_of_candidates = cugraph::collect_values_for_keys(handle, + target_candidate_map.view(), + candidates.begin(), + candidates.end(), + vertex_to_gpu_id_op); + } else { + candidates_of_candidates.resize(candidates.size(), handle.get_stream()); + + target_candidate_map.view().find(candidates.begin(), + candidates.end(), + candidates_of_candidates.begin(), + handle.get_stream()); + } + + // + // Mask out neighborhood of matched vertices + // + + rmm::device_uvector is_vertex_matched = rmm::device_uvector( + current_graph_view.local_vertex_partition_range_size(), handle.get_stream()); + thrust::fill( + handle.get_thrust_policy(), is_vertex_matched.begin(), is_vertex_matched.end(), bool{false}); + + thrust::for_each( + handle.get_thrust_policy(), + thrust::make_zip_iterator(thrust::make_tuple(candidates_of_candidates.begin(), + targets.begin(), + candidates.begin(), + offers_from_candidates.begin())), + thrust::make_zip_iterator(thrust::make_tuple(candidates_of_candidates.end(), + targets.end(), + candidates.end(), + offers_from_candidates.end())), + [partners = partners.begin(), + offers_from_partners = offers_from_partners.begin(), + is_vertex_matched = + raft::device_span(is_vertex_matched.data(), is_vertex_matched.size()), + v_first = + current_graph_view.local_vertex_partition_range_first()] __device__(auto msrc_tgt) { + auto candidate_of_candidate = thrust::get<0>(msrc_tgt); + auto tgt = thrust::get<1>(msrc_tgt); + auto candiate = thrust::get<2>(msrc_tgt); + auto offer_value = thrust::get<3>(msrc_tgt); + + if (candidate_of_candidate != invalid_partner && candidate_of_candidate == tgt) { + auto tgt_offset = tgt - v_first; + is_vertex_matched[tgt_offset] = true; + partners[tgt_offset] = candiate; + offers_from_partners[tgt_offset] = offer_value; + } + }); + + if (current_graph_view.compute_number_of_edges(handle) == 0) { break; } + + if constexpr (graph_view_t::is_multi_gpu) { + cugraph::update_edge_src_property( + handle, current_graph_view, is_vertex_matched.begin(), src_match_flags); + cugraph::update_edge_dst_property( + handle, current_graph_view, is_vertex_matched.begin(), dst_match_flags); + } + + if (loop_counter % 2 == 0) { + if constexpr (graph_view_t::is_multi_gpu) { + cugraph::transform_e( + handle, + current_graph_view, + src_match_flags.view(), + dst_match_flags.view(), + cugraph::edge_dummy_property_t{}.view(), + [] __device__( + auto src, auto dst, auto is_src_matched, auto is_dst_matched, thrust::nullopt_t) { + return !((is_src_matched == true) || (is_dst_matched == true)); + }, + edge_masks_odd.mutable_view()); + } else { + cugraph::transform_e( + handle, + current_graph_view, + detail::edge_major_property_view_t(is_vertex_matched.begin()), + detail::edge_minor_property_view_t(is_vertex_matched.begin(), + vertex_t{0}), + cugraph::edge_dummy_property_t{}.view(), + [] __device__( + auto src, auto dst, auto is_src_matched, auto is_dst_matched, thrust::nullopt_t) { + return !((is_src_matched == true) || (is_dst_matched == true)); + }, + edge_masks_odd.mutable_view()); + } + + if (current_graph_view.has_edge_mask()) current_graph_view.clear_edge_mask(); + cugraph::fill_edge_property(handle, current_graph_view, bool{false}, edge_masks_even); + current_graph_view.attach_edge_mask(edge_masks_odd.view()); + } else { + if constexpr (graph_view_t::is_multi_gpu) { + cugraph::transform_e( + handle, + current_graph_view, + src_match_flags.view(), + dst_match_flags.view(), + cugraph::edge_dummy_property_t{}.view(), + [] __device__( + auto src, auto dst, auto is_src_matched, auto is_dst_matched, thrust::nullopt_t) { + return !((is_src_matched == true) || (is_dst_matched == true)); + }, + edge_masks_even.mutable_view()); + } else { + cugraph::transform_e( + handle, + current_graph_view, + detail::edge_major_property_view_t(is_vertex_matched.begin()), + detail::edge_minor_property_view_t(is_vertex_matched.begin(), + vertex_t{0}), + cugraph::edge_dummy_property_t{}.view(), + [] __device__( + auto src, auto dst, auto is_src_matched, auto is_dst_matched, thrust::nullopt_t) { + return !((is_src_matched == true) || (is_dst_matched == true)); + }, + edge_masks_even.mutable_view()); + } + + if (current_graph_view.has_edge_mask()) current_graph_view.clear_edge_mask(); + cugraph::fill_edge_property(handle, current_graph_view, bool{false}, edge_masks_odd); + current_graph_view.attach_edge_mask(edge_masks_even.view()); + } + + loop_counter++; + } + + weight_t sum_matched_edge_weights = thrust::reduce( + handle.get_thrust_policy(), offers_from_partners.begin(), offers_from_partners.end()); + + if constexpr (graph_view_t::is_multi_gpu) { + sum_matched_edge_weights = host_scalar_allreduce( + handle.get_comms(), sum_matched_edge_weights, raft::comms::op_t::SUM, handle.get_stream()); + } + + return std::make_tuple(std::move(partners), sum_matched_edge_weights / 2.0); +} +} // namespace detail + +template +std::tuple, weight_t> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view) +{ + return detail::approximate_weighted_matching(handle, graph_view, edge_weight_view); +} + +} // namespace cugraph diff --git a/cpp/src/community/approx_weighted_matching_mg.cu b/cpp/src/community/approx_weighted_matching_mg.cu new file mode 100644 index 00000000000..41d6c3d97e0 --- /dev/null +++ b/cpp/src/community/approx_weighted_matching_mg.cu @@ -0,0 +1,50 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include "approx_weighted_matching_impl.cuh" + +namespace cugraph { + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +} // namespace cugraph diff --git a/cpp/src/community/approx_weighted_matching_sg.cu b/cpp/src/community/approx_weighted_matching_sg.cu new file mode 100644 index 00000000000..418a43d51ae --- /dev/null +++ b/cpp/src/community/approx_weighted_matching_sg.cu @@ -0,0 +1,50 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include "approx_weighted_matching_impl.cuh" + +namespace cugraph { + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, float> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +template std::tuple, double> approximate_weighted_matching( + raft::handle_t const& handle, + graph_view_t const& graph_view, + edge_property_view_t edge_weight_view); + +} // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 19097add541..ced3b7bedb1 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -309,6 +309,10 @@ ConfigureTest(LOUVAIN_TEST community/louvain_test.cpp) # - LEIDEN tests ---------------------------------------------------------------------------------- ConfigureTest(LEIDEN_TEST community/leiden_test.cpp) +################################################################################################### +# - WEIGHTED MATCHING tests ---------------------------------------------------------------------------------- +ConfigureTest(WEIGHTED_MATCHING_TEST community/weighted_matching_test.cpp) + ################################################################################################### # - Legacy ECG tests ------------------------------------------------------------------------------------- ConfigureTest(LEGACY_ECG_TEST community/legacy_ecg_test.cpp) @@ -570,6 +574,10 @@ if(BUILD_CUGRAPH_MG_TESTS) # - MG LEIDEN tests -------------------------------------------------------------------------- ConfigureTestMG(MG_LEIDEN_TEST community/mg_leiden_test.cpp) + ############################################################################################### + # - MG WEIGHTED MATCHING tests -------------------------------------------------------------------------- + ConfigureTestMG(MG_WEIGHTED_MATCHING_TEST community/mg_weighted_matching_test.cpp) + ############################################################################################### # - MG ECG tests -------------------------------------------------------------------------- ConfigureTestMG(MG_ECG_TEST community/mg_ecg_test.cpp) diff --git a/cpp/tests/community/mg_weighted_matching_test.cpp b/cpp/tests/community/mg_weighted_matching_test.cpp new file mode 100644 index 00000000000..21963922ab1 --- /dev/null +++ b/cpp/tests/community/mg_weighted_matching_test.cpp @@ -0,0 +1,232 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governin_from_mtxg permissions and + * limitations under the License. + */ + +#include "utilities/base_fixture.hpp" +#include "utilities/conversion_utilities.hpp" +#include "utilities/property_generator_utilities.hpp" +#include "utilities/test_graphs.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +struct WeightedMatching_UseCase { + bool edge_masking{false}; + bool check_correctness{true}; +}; + +template +class Tests_MGWeightedMatching + : public ::testing::TestWithParam> { + public: + Tests_MGWeightedMatching() {} + + static void SetUpTestCase() { handle_ = cugraph::test::initialize_mg_handle(); } + static void TearDownTestCase() { handle_.reset(); } + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [weighted_matching_usecase, input_usecase] = param; + + HighResTimer hr_timer{}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + handle_->get_comms().barrier(); + hr_timer.start("MG Construct graph"); + } + + constexpr bool multi_gpu = true; + + bool test_weighted = true; + bool renumber = true; + bool drop_self_loops = false; + bool drop_multi_edges = false; + + auto [mg_graph, mg_edge_weights, mg_renumber_map] = + cugraph::test::construct_graph( + *handle_, input_usecase, test_weighted, renumber, drop_self_loops, drop_multi_edges); + + std::tie(mg_graph, mg_edge_weights, mg_renumber_map) = cugraph::symmetrize_graph( + *handle_, + std::move(mg_graph), + std::move(mg_edge_weights), + mg_renumber_map ? std::optional>(std::move(*mg_renumber_map)) + : std::nullopt, + false); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + handle_->get_comms().barrier(); + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + auto mg_graph_view = mg_graph.view(); + auto mg_edge_weight_view = + mg_edge_weights ? std::make_optional((*mg_edge_weights).view()) : std::nullopt; + + std::optional> edge_mask{std::nullopt}; + if (weighted_matching_usecase.edge_masking) { + edge_mask = cugraph::test::generate::edge_property( + *handle_, mg_graph_view, 2); + mg_graph_view.attach_edge_mask((*edge_mask).view()); + } + + rmm::device_uvector mg_partners(0, handle_->get_stream()); + weight_t mg_matching_weights; + + std::forward_as_tuple(mg_partners, mg_matching_weights) = + cugraph::approximate_weighted_matching( + *handle_, mg_graph_view, (*mg_edge_weights).view()); + + if (weighted_matching_usecase.check_correctness) { + auto h_mg_partners = cugraph::test::to_host(*handle_, mg_partners); + + auto constexpr invalid_partner = cugraph::invalid_vertex_id::value; + + rmm::device_uvector mg_aggregate_partners(0, handle_->get_stream()); + std::tie(std::ignore, mg_aggregate_partners) = + cugraph::test::mg_vertex_property_values_to_sg_vertex_property_values( + *handle_, + std::optional>{std::nullopt}, + mg_graph_view.local_vertex_partition_range(), + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + raft::device_span(mg_partners.data(), mg_partners.size())); + + cugraph::graph_t sg_graph(*handle_); + std::optional< + cugraph::edge_property_t, weight_t>> + sg_edge_weights{std::nullopt}; + std::tie(sg_graph, sg_edge_weights, std::ignore) = cugraph::test::mg_graph_to_sg_graph( + *handle_, + mg_graph_view, + mg_edge_weight_view, + std::optional>(std::nullopt), + false); + + if (handle_->get_comms().get_rank() == 0) { + auto sg_graph_view = sg_graph.view(); + + rmm::device_uvector sg_partners(0, handle_->get_stream()); + weight_t sg_matching_weights; + + std::forward_as_tuple(sg_partners, sg_matching_weights) = + cugraph::approximate_weighted_matching( + *handle_, sg_graph_view, (*sg_edge_weights).view()); + auto h_sg_partners = cugraph::test::to_host(*handle_, sg_partners); + auto h_mg_aggregate_partners = cugraph::test::to_host(*handle_, mg_aggregate_partners); + + ASSERT_FLOAT_EQ(mg_matching_weights, sg_matching_weights) + << "SG and MG matching weights are different"; + ASSERT_TRUE( + std::equal(h_sg_partners.begin(), h_sg_partners.end(), h_mg_aggregate_partners.begin())); + } + } + } + + private: + static std::unique_ptr handle_; +}; + +template +std::unique_ptr Tests_MGWeightedMatching::handle_ = nullptr; + +using Tests_MGWeightedMatching_File = Tests_MGWeightedMatching; +using Tests_MGWeightedMatching_Rmat = Tests_MGWeightedMatching; + +TEST_P(Tests_MGWeightedMatching_File, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGWeightedMatching_File, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGWeightedMatching_File, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGWeightedMatching_Rmat, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGWeightedMatching_Rmat, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGWeightedMatching_Rmat, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_MGWeightedMatching_File, + ::testing::Combine(::testing::Values(WeightedMatching_UseCase{false}, + WeightedMatching_UseCase{true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P(rmat_small_test, + Tests_MGWeightedMatching_Rmat, + ::testing::Combine(::testing::Values(WeightedMatching_UseCase{false}, + WeightedMatching_UseCase{true}), + ::testing::Values(cugraph::test::Rmat_Usecase( + 3, 2, 0.57, 0.19, 0.19, 0, true, false)))); + +INSTANTIATE_TEST_SUITE_P( + rmat_benchmark_test, /* note that scale & edge factor can be overridden in benchmarking (with + --gtest_filter to select only the rmat_benchmark_test with a specific + vertex & edge type combination) by command line arguments and do not + include more than one Rmat_Usecase that differ only in scale or edge + factor (to avoid running same benchmarks more than once) */ + Tests_MGWeightedMatching_Rmat, + ::testing::Combine( + ::testing::Values(WeightedMatching_UseCase{false, false}, + WeightedMatching_UseCase{true, false}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, true, false)))); + +CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/community/weighted_matching_test.cpp b/cpp/tests/community/weighted_matching_test.cpp new file mode 100644 index 00000000000..436273c3be3 --- /dev/null +++ b/cpp/tests/community/weighted_matching_test.cpp @@ -0,0 +1,182 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governin_from_mtxg permissions and + * limitations under the License. + */ + +#include "utilities/base_fixture.hpp" +#include "utilities/conversion_utilities.hpp" +#include "utilities/property_generator_utilities.hpp" +#include "utilities/test_graphs.hpp" + +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +struct WeightedMatching_UseCase { + bool edge_masking{false}; + bool check_correctness{true}; +}; + +template +class Tests_SGWeightedMatching + : public ::testing::TestWithParam> { + public: + Tests_SGWeightedMatching() {} + + static void SetUpTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [weighted_matching_usecase, input_usecase] = param; + + raft::handle_t handle{}; + HighResTimer hr_timer{}; + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + hr_timer.start("Construct graph"); + } + + constexpr bool multi_gpu = false; + + bool test_weighted = true; + bool renumber = true; + bool drop_self_loops = false; + bool drop_multi_edges = false; + + auto [sg_graph, sg_edge_weights, sg_renumber_map] = + cugraph::test::construct_graph( + handle, input_usecase, test_weighted, renumber, drop_self_loops, drop_multi_edges); + + std::tie(sg_graph, sg_edge_weights, sg_renumber_map) = cugraph::symmetrize_graph( + handle, std::move(sg_graph), std::move(sg_edge_weights), std::move(sg_renumber_map), false); + + if (cugraph::test::g_perf) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + hr_timer.stop(); + hr_timer.display_and_clear(std::cout); + } + + auto sg_graph_view = sg_graph.view(); + auto sg_edge_weight_view = + sg_edge_weights ? std::make_optional((*sg_edge_weights).view()) : std::nullopt; + + std::optional> edge_mask{std::nullopt}; + if (weighted_matching_usecase.edge_masking) { + edge_mask = cugraph::test::generate::edge_property( + handle, sg_graph_view, 2); + sg_graph_view.attach_edge_mask((*edge_mask).view()); + } + + rmm::device_uvector d_partners(0, handle.get_stream()); + weight_t total_matching_weights; + + std::forward_as_tuple(d_partners, total_matching_weights) = + cugraph::approximate_weighted_matching( + handle, sg_graph_view, (*sg_edge_weights).view()); + + if (weighted_matching_usecase.check_correctness) { + auto h_partners = cugraph::test::to_host(handle, d_partners); + auto constexpr invalid_partner = cugraph::invalid_vertex_id::value; + + std::for_each(h_partners.begin(), h_partners.end(), [&invalid_partner, h_partners](auto& v) { + if (v != invalid_partner) ASSERT_TRUE(h_partners[h_partners[v]] == v); + }); + } + } +}; + +using Tests_SGWeightedMatching_File = Tests_SGWeightedMatching; +using Tests_SGWeightedMatching_Rmat = Tests_SGWeightedMatching; + +TEST_P(Tests_SGWeightedMatching_File, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGWeightedMatching_File, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGWeightedMatching_File, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGWeightedMatching_Rmat, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGWeightedMatching_Rmat, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGWeightedMatching_Rmat, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_SGWeightedMatching_File, + ::testing::Combine(::testing::Values(WeightedMatching_UseCase{false}, + WeightedMatching_UseCase{true}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P(rmat_small_test, + Tests_SGWeightedMatching_Rmat, + ::testing::Combine(::testing::Values(WeightedMatching_UseCase{false}, + WeightedMatching_UseCase{true}), + ::testing::Values(cugraph::test::Rmat_Usecase( + 3, 3, 0.57, 0.19, 0.19, 0, true, false)))); + +INSTANTIATE_TEST_SUITE_P( + rmat_benchmark_test, /* note that scale & edge factor can be overridden in benchmarking (with + --gtest_filter to select only the rmat_benchmark_test with a specific + vertex & edge type combination) by command line arguments and do not + include more than one Rmat_Usecase that differ only in scale or edge + factor (to avoid running same benchmarks more than once) */ + Tests_SGWeightedMatching_Rmat, + ::testing::Combine( + ::testing::Values(WeightedMatching_UseCase{false, false}, + WeightedMatching_UseCase{true, false}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); + +CUGRAPH_TEST_PROGRAM_MAIN()