From 0e0b36371bfc5d3055007d596e383d5e12ddb457 Mon Sep 17 00:00:00 2001 From: Anthony Mahanna <43019056+aMahanna@users.noreply.github.com> Date: Tue, 12 Mar 2024 12:40:08 -0400 Subject: [PATCH 1/3] [docs] fix GNN typo (#4196) Authors: - Anthony Mahanna (https://github.com/aMahanna) - Alex Barghi (https://github.com/alexbarghi-nv) - Ralph Liu (https://github.com/nv-rliu) - Brad Rees (https://github.com/BradReesWork) Approvers: - Don Acosta (https://github.com/acostadon) URL: https://github.com/rapidsai/cugraph/pull/4196 From e4b58ec04e4add07578ce779977835cfe28f516a Mon Sep 17 00:00:00 2001 From: Naim <110031745+naimnv@users.noreply.github.com> Date: Tue, 12 Mar 2024 19:00:31 +0100 Subject: [PATCH 2/3] Make external vertex and edge shuffling function public (#4227) Make vertex and edge shuffling function public Authors: - Naim (https://github.com/naimnv) - Ralph Liu (https://github.com/nv-rliu) Approvers: - Chuck Hastings (https://github.com/ChuckHastings) - Seunghwa Kang (https://github.com/seunghwak) URL: https://github.com/rapidsai/cugraph/pull/4227 --- cpp/CMakeLists.txt | 4 +- cpp/include/cugraph/graph_functions.hpp | 63 +++++++++++ .../shuffle_vertex_pairs.cu | 102 ++++++++++++++++++ .../{detail => utilities}/shuffle_vertices.cu | 70 ++++++++++++ 4 files changed, 237 insertions(+), 2 deletions(-) rename cpp/src/{detail => utilities}/shuffle_vertex_pairs.cu (80%) rename cpp/src/{detail => utilities}/shuffle_vertices.cu (73%) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index a3392627fb8..3131404712f 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -186,9 +186,9 @@ endif() # which should give us a better parallel schedule. set(CUGRAPH_SOURCES - src/detail/shuffle_vertices.cu + src/utilities/shuffle_vertices.cu src/detail/permute_range.cu - src/detail/shuffle_vertex_pairs.cu + src/utilities/shuffle_vertex_pairs.cu src/detail/collect_local_vertex_values.cu src/detail/groupby_and_count.cu src/detail/collect_comm_wrapper.cu diff --git a/cpp/include/cugraph/graph_functions.hpp b/cpp/include/cugraph/graph_functions.hpp index 90425f86bef..6d4470e8251 100644 --- a/cpp/include/cugraph/graph_functions.hpp +++ b/cpp/include/cugraph/graph_functions.hpp @@ -1052,4 +1052,67 @@ remove_multi_edges(raft::handle_t const& handle, std::optional>&& edgelist_edge_types, bool keep_min_value_edge = false); +/** + * @brief Shuffle external vertex ids to the proper GPU. + * + * @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 vertices List of vertex ids + * @return Vector of vertex ids mapped to this GPU. + */ +template +rmm::device_uvector shuffle_external_vertices(raft::handle_t const& handle, + rmm::device_uvector&& vertices); + +/** + * @brief Shuffle external vertex ids and values to the proper GPU. + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam value_t Type of values. currently supported types are int32_t, + * int64_t, size_t, float and double. + * + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param vertices List of vertex ids + * @param values List of values + * @return Tuple of vectors storing vertex ids and values mapped to this GPU. + */ +template +std::tuple, rmm::device_uvector> +shuffle_external_vertex_value_pairs(raft::handle_t const& handle, + rmm::device_uvector&& vertices, + rmm::device_uvector&& values); + +/** + * @brief Shuffle external edges to the proper GPU. + * + * @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. + * + * @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 List of source vertex ids + * @param edge_dsts List of destination vertex ids + * @param edge_weights Optional list of edge weights + * @param edge_ids Optional list of edge ids + * @param edge_types Optional list of edge types + * @return Tuple of vectors storing edge sources, destinations, optional weights, + * optional edge ids, optional edge types mapped to this GPU. + */ +template +std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +shuffle_external_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); + } // namespace cugraph diff --git a/cpp/src/detail/shuffle_vertex_pairs.cu b/cpp/src/utilities/shuffle_vertex_pairs.cu similarity index 80% rename from cpp/src/detail/shuffle_vertex_pairs.cu rename to cpp/src/utilities/shuffle_vertex_pairs.cu index 33a7834f5ff..b473796aa9d 100644 --- a/cpp/src/detail/shuffle_vertex_pairs.cu +++ b/cpp/src/utilities/shuffle_vertex_pairs.cu @@ -519,4 +519,106 @@ shuffle_int_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( std::vector const& vertex_partition_range_lasts); } // namespace detail + +template +std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +shuffle_external_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) +{ + auto& comm = handle.get_comms(); + auto const comm_size = comm.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(); + + return detail::shuffle_ext_vertex_pairs_with_values_to_local_gpu_by_edge_partitioning( + handle, + std::move(edge_srcs), + std::move(edge_dsts), + std::move(edge_weights), + std::move(edge_ids), + std::move(edge_types)); +} + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +shuffle_external_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); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +shuffle_external_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); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +shuffle_external_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); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +shuffle_external_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); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +shuffle_external_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); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>> +shuffle_external_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); + } // namespace cugraph diff --git a/cpp/src/detail/shuffle_vertices.cu b/cpp/src/utilities/shuffle_vertices.cu similarity index 73% rename from cpp/src/detail/shuffle_vertices.cu rename to cpp/src/utilities/shuffle_vertices.cu index be6875f1073..b396201f509 100644 --- a/cpp/src/detail/shuffle_vertices.cu +++ b/cpp/src/utilities/shuffle_vertices.cu @@ -249,4 +249,74 @@ shuffle_ext_vertex_value_pairs_to_local_gpu_by_vertex_partitioning( rmm::device_uvector&& values); } // namespace detail + +template +std::tuple, rmm::device_uvector> +shuffle_external_vertex_value_pairs(raft::handle_t const& handle, + rmm::device_uvector&& vertices, + rmm::device_uvector&& values) +{ + return detail::shuffle_ext_vertex_value_pairs_to_local_gpu_by_vertex_partitioning( + handle, std::move(vertices), std::move(values)); +} + +template std::tuple, rmm::device_uvector> +shuffle_external_vertex_value_pairs(raft::handle_t const& handle, + rmm::device_uvector&& vertices, + rmm::device_uvector&& values); + +template std::tuple, rmm::device_uvector> +shuffle_external_vertex_value_pairs(raft::handle_t const& handle, + rmm::device_uvector&& vertices, + rmm::device_uvector&& values); + +template std::tuple, rmm::device_uvector> +shuffle_external_vertex_value_pairs(raft::handle_t const& handle, + rmm::device_uvector&& vertices, + rmm::device_uvector&& values); + +template std::tuple, rmm::device_uvector> +shuffle_external_vertex_value_pairs(raft::handle_t const& handle, + rmm::device_uvector&& vertices, + rmm::device_uvector&& values); + +template std::tuple, rmm::device_uvector> +shuffle_external_vertex_value_pairs(raft::handle_t const& handle, + rmm::device_uvector&& vertices, + rmm::device_uvector&& values); + +template std::tuple, rmm::device_uvector> +shuffle_external_vertex_value_pairs(raft::handle_t const& handle, + rmm::device_uvector&& vertices, + rmm::device_uvector&& values); + +template std::tuple, rmm::device_uvector> +shuffle_external_vertex_value_pairs(raft::handle_t const& handle, + rmm::device_uvector&& vertices, + rmm::device_uvector&& values); + +template std::tuple, rmm::device_uvector> +shuffle_external_vertex_value_pairs(raft::handle_t const& handle, + rmm::device_uvector&& vertices, + rmm::device_uvector&& values); + +template std::tuple, rmm::device_uvector> +shuffle_external_vertex_value_pairs(raft::handle_t const& handle, + rmm::device_uvector&& vertices, + rmm::device_uvector&& values); + +template +rmm::device_uvector shuffle_external_vertices(raft::handle_t const& handle, + rmm::device_uvector&& vertices) +{ + return detail::shuffle_ext_vertices_to_local_gpu_by_vertex_partitioning(handle, + std::move(vertices)); +} + +template rmm::device_uvector shuffle_external_vertices( + raft::handle_t const& handle, rmm::device_uvector&& d_vertices); + +template rmm::device_uvector shuffle_external_vertices( + raft::handle_t const& handle, rmm::device_uvector&& d_vertices); + } // namespace cugraph From a4eab99f8f0ff50aab980dfab1fffa68e8983942 Mon Sep 17 00:00:00 2001 From: Naim <110031745+naimnv@users.noreply.github.com> Date: Tue, 12 Mar 2024 19:01:20 +0100 Subject: [PATCH 3/3] Graph coloring and MIS (#4211) This PR implements - MNMG Graph Coloring - MNMG MIS for general graphs Closes #4230 Closes #4231 Authors: - Naim (https://github.com/naimnv) Approvers: - Seunghwa Kang (https://github.com/seunghwak) - Chuck Hastings (https://github.com/ChuckHastings) URL: https://github.com/rapidsai/cugraph/pull/4211 --- cpp/CMakeLists.txt | 8 +- cpp/include/cugraph/algorithms.hpp | 30 +- .../detail/maximal_independent_moves.cuh | 313 ++++++++++++++++++ .../maximal_independent_moves.hpp} | 6 +- .../detail/maximal_independent_moves_mg.cu | 38 +++ .../detail/maximal_independent_moves_sg.cu | 37 +++ cpp/src/community/detail/refine_impl.cuh | 6 +- .../detail => components}/mis_impl.cuh | 37 ++- .../detail => components}/mis_mg.cu | 8 +- .../detail => components}/mis_sg.cu | 7 +- cpp/src/components/vertex_coloring_impl.cuh | 150 +++++++++ cpp/src/components/vertex_coloring_mg.cu | 35 ++ cpp/src/components/vertex_coloring_sg.cu | 35 ++ cpp/tests/CMakeLists.txt | 20 +- .../{community => components}/mg_mis_test.cu | 19 +- .../components/mg_vertex_coloring_test.cu | 275 +++++++++++++++ cpp/tests/components/mis_test.cu | 240 ++++++++++++++ cpp/tests/components/vertex_coloring_test.cu | 244 ++++++++++++++ 18 files changed, 1468 insertions(+), 40 deletions(-) create mode 100644 cpp/src/community/detail/maximal_independent_moves.cuh rename cpp/src/community/{mis.hpp => detail/maximal_independent_moves.hpp} (91%) create mode 100644 cpp/src/community/detail/maximal_independent_moves_mg.cu create mode 100644 cpp/src/community/detail/maximal_independent_moves_sg.cu rename cpp/src/{community/detail => components}/mis_impl.cuh (90%) rename cpp/src/{community/detail => components}/mis_mg.cu (83%) rename cpp/src/{community/detail => components}/mis_sg.cu (82%) create mode 100644 cpp/src/components/vertex_coloring_impl.cuh create mode 100644 cpp/src/components/vertex_coloring_mg.cu create mode 100644 cpp/src/components/vertex_coloring_sg.cu rename cpp/tests/{community => components}/mg_mis_test.cu (95%) create mode 100644 cpp/tests/components/mg_vertex_coloring_test.cu create mode 100644 cpp/tests/components/mis_test.cu create mode 100644 cpp/tests/components/vertex_coloring_test.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 3131404712f..b12403710ab 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -197,8 +197,8 @@ set(CUGRAPH_SOURCES src/community/detail/common_methods_sg.cu src/community/detail/refine_sg.cu src/community/detail/refine_mg.cu - src/community/detail/mis_sg.cu - src/community/detail/mis_mg.cu + src/community/detail/maximal_independent_moves_sg.cu + src/community/detail/maximal_independent_moves_mg.cu src/detail/utility_wrappers.cu src/structure/graph_view_mg.cu src/structure/remove_self_loops.cu @@ -295,6 +295,10 @@ set(CUGRAPH_SOURCES src/tree/legacy/mst.cu src/components/weakly_connected_components_sg.cu src/components/weakly_connected_components_mg.cu + src/components/mis_sg.cu + src/components/mis_mg.cu + src/components/vertex_coloring_sg.cu + src/components/vertex_coloring_mg.cu src/structure/create_graph_from_edgelist_sg.cu src/structure/create_graph_from_edgelist_mg.cu src/structure/symmetrize_edgelist_sg.cu diff --git a/cpp/include/cugraph/algorithms.hpp b/cpp/include/cugraph/algorithms.hpp index 5c29604a5a7..1471d340cec 100644 --- a/cpp/include/cugraph/algorithms.hpp +++ b/cpp/include/cugraph/algorithms.hpp @@ -2340,15 +2340,41 @@ std::tuple, rmm::device_uvector> k_hop_nbr * handles to various CUDA libraries) to run graph algorithms. * @param graph_view Graph view object. * @param rng_state The RngState instance holding pseudo-random number generator state. - * @return A device vector containing vertices found in the maximal independent set + * @return A device vector containing vertices in the maximal independent set. */ - template rmm::device_uvector maximal_independent_set( raft::handle_t const& handle, graph_view_t const& graph_view, raft::random::RngState& rng_state); +/* + * @brief Find a Greedy Vertex Coloring + * + * A vertex coloring is an assignment of colors or labels to each vertex of a graph so that + * no two adjacent vertices have the same color or label. Finding the minimum number of colors + * needed to color the vertices of a graph is an NP-hard problem and therefore for practical + * use cases greedy coloring is used. Here we provide an implementation of greedy vertex + * coloring based on maximal independent set. + * See + * https://research.nvidia.com/sites/default/files/pubs/2015-05_Parallel-Graph-Coloring/nvr-2015-001.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 handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param graph_view Graph view object. + * @param rng_state The RngState instance holding pseudo-random number generator state. + * @return A device vector containing color for each vertex. + */ +template +rmm::device_uvector vertex_coloring( + raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::random::RngState& rng_state); + } // namespace cugraph /** diff --git a/cpp/src/community/detail/maximal_independent_moves.cuh b/cpp/src/community/detail/maximal_independent_moves.cuh new file mode 100644 index 00000000000..82d20a04203 --- /dev/null +++ b/cpp/src/community/detail/maximal_independent_moves.cuh @@ -0,0 +1,313 @@ + +/* + * Copyright (c) 2023-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 "maximal_independent_moves.hpp" +#include "prims/fill_edge_src_dst_property.cuh" +#include "prims/per_v_transform_reduce_incoming_outgoing_e.cuh" +#include "prims/update_edge_src_dst_property.cuh" + +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +namespace cugraph { + +namespace detail { + +template +rmm::device_uvector maximal_independent_moves( + raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::random::RngState& rng_state) +{ + using GraphViewType = cugraph::graph_view_t; + + vertex_t local_vtx_partitoin_size = graph_view.local_vertex_partition_range_size(); + + rmm::device_uvector remaining_vertices(local_vtx_partitoin_size, handle.get_stream()); + + auto vertex_begin = + thrust::make_counting_iterator(graph_view.local_vertex_partition_range_first()); + auto vertex_end = thrust::make_counting_iterator(graph_view.local_vertex_partition_range_last()); + + // Compute out-degree + auto out_degrees = graph_view.compute_out_degrees(handle); + + // Only vertices with non-zero out-degree are possible can move + remaining_vertices.resize( + thrust::distance(remaining_vertices.begin(), + thrust::copy_if(handle.get_thrust_policy(), + vertex_begin, + vertex_end, + out_degrees.begin(), + remaining_vertices.begin(), + [] __device__(auto deg) { return deg > 0; })), + handle.get_stream()); + + // Set ID of each vertex as its rank + rmm::device_uvector ranks(local_vtx_partitoin_size, handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), vertex_begin, vertex_end, ranks.begin()); + + // Set ranks of zero out-degree vetices to std::numeric_limits::lowest() + thrust::transform_if(handle.get_thrust_policy(), + out_degrees.begin(), + out_degrees.end(), + ranks.begin(), + cuda::proclaim_return_type( + [] __device__(auto) { return std::numeric_limits::lowest(); }), + [] __device__(auto deg) { return deg == 0; }); + + out_degrees.resize(0, handle.get_stream()); + out_degrees.shrink_to_fit(handle.get_stream()); + + size_t loop_counter = 0; + while (true) { + loop_counter++; + + // Copy ranks into temporary vector to begin with + + rmm::device_uvector temporary_ranks(local_vtx_partitoin_size, handle.get_stream()); + thrust::copy(handle.get_thrust_policy(), ranks.begin(), ranks.end(), temporary_ranks.begin()); + + // Select a random set of candidate vertices + + vertex_t nr_remaining_vertices_to_check = remaining_vertices.size(); + if (multi_gpu) { + nr_remaining_vertices_to_check = host_scalar_allreduce(handle.get_comms(), + nr_remaining_vertices_to_check, + raft::comms::op_t::SUM, + handle.get_stream()); + } + + vertex_t nr_candidates = (nr_remaining_vertices_to_check < 1024) + ? nr_remaining_vertices_to_check + : std::min(static_cast((0.50 + 0.25 * loop_counter) * + nr_remaining_vertices_to_check), + nr_remaining_vertices_to_check); + + // FIXME: Can we improve performance here? + // FIXME: if(nr_remaining_vertices_to_check < 1024), may avoid calling select_random_vertices + auto d_sampled_vertices = + cugraph::select_random_vertices(handle, + graph_view, + std::make_optional(raft::device_span{ + remaining_vertices.data(), remaining_vertices.size()}), + rng_state, + nr_candidates, + false, + true); + + rmm::device_uvector non_candidate_vertices( + remaining_vertices.size() - d_sampled_vertices.size(), handle.get_stream()); + + thrust::set_difference(handle.get_thrust_policy(), + remaining_vertices.begin(), + remaining_vertices.end(), + d_sampled_vertices.begin(), + d_sampled_vertices.end(), + non_candidate_vertices.begin()); + + // Set temporary ranks of non-candidate vertices to std::numeric_limits::lowest() + thrust::for_each( + handle.get_thrust_policy(), + non_candidate_vertices.begin(), + non_candidate_vertices.end(), + [temporary_ranks = + raft::device_span(temporary_ranks.data(), temporary_ranks.size()), + v_first = graph_view.local_vertex_partition_range_first()] __device__(auto v) { + // + // if rank of a non-candidate vertex is not std::numeric_limits::max() (i.e. the + // vertex is not already in MIS), set it to std::numeric_limits::lowest() + // + auto v_offset = v - v_first; + if (temporary_ranks[v_offset] < std::numeric_limits::max()) { + temporary_ranks[v_offset] = std::numeric_limits::lowest(); + } + }); + + // Caches for ranks + edge_src_property_t src_rank_cache(handle); + edge_dst_property_t dst_rank_cache(handle); + + // Update rank caches with temporary ranks + if constexpr (multi_gpu) { + src_rank_cache = edge_src_property_t(handle, graph_view); + dst_rank_cache = edge_dst_property_t(handle, graph_view); + update_edge_src_property(handle, graph_view, temporary_ranks.begin(), src_rank_cache); + update_edge_dst_property(handle, graph_view, temporary_ranks.begin(), dst_rank_cache); + } + + // + // Find maximum rank outgoing neighbor for each vertex + // + + rmm::device_uvector max_outgoing_ranks(local_vtx_partitoin_size, handle.get_stream()); + + per_v_transform_reduce_outgoing_e( + handle, + graph_view, + multi_gpu + ? src_rank_cache.view() + : detail::edge_major_property_view_t(temporary_ranks.data()), + multi_gpu ? dst_rank_cache.view() + : detail::edge_minor_property_view_t( + temporary_ranks.data(), vertex_t{0}), + edge_dummy_property_t{}.view(), + [] __device__(auto src, auto dst, auto src_rank, auto dst_rank, auto wt) { return dst_rank; }, + std::numeric_limits::lowest(), + cugraph::reduce_op::maximum{}, + max_outgoing_ranks.begin()); + + // + // Find maximum rank incoming neighbor for each vertex + // + + rmm::device_uvector max_incoming_ranks(local_vtx_partitoin_size, handle.get_stream()); + + per_v_transform_reduce_incoming_e( + handle, + graph_view, + multi_gpu + ? src_rank_cache.view() + : detail::edge_major_property_view_t(temporary_ranks.data()), + multi_gpu ? dst_rank_cache.view() + : detail::edge_minor_property_view_t( + temporary_ranks.data(), vertex_t{0}), + edge_dummy_property_t{}.view(), + [] __device__(auto src, auto dst, auto src_rank, auto dst_rank, auto wt) { return src_rank; }, + std::numeric_limits::lowest(), + cugraph::reduce_op::maximum{}, + max_incoming_ranks.begin()); + + temporary_ranks.resize(0, handle.get_stream()); + temporary_ranks.shrink_to_fit(handle.get_stream()); + + // + // Compute max of outgoing and incoming neighbors + // + thrust::transform(handle.get_thrust_policy(), + max_incoming_ranks.begin(), + max_incoming_ranks.end(), + max_outgoing_ranks.begin(), + max_outgoing_ranks.begin(), + thrust::maximum()); + + max_incoming_ranks.resize(0, handle.get_stream()); + max_incoming_ranks.shrink_to_fit(handle.get_stream()); + + // + // If the max neighbor of a vertex is already in MIS (i.e. has rank + // std::numeric_limits::max()), discard it, otherwise, + // include the vertex if it has larger rank than its maximum rank neighbor + // + auto last = thrust::remove_if( + handle.get_thrust_policy(), + d_sampled_vertices.begin(), + d_sampled_vertices.end(), + [max_rank_neighbor_first = max_outgoing_ranks.begin(), + ranks = raft::device_span(ranks.data(), ranks.size()), + v_first = graph_view.local_vertex_partition_range_first()] __device__(auto v) { + auto v_offset = v - v_first; + auto max_neighbor_rank = *(max_rank_neighbor_first + v_offset); + auto rank_of_v = ranks[v_offset]; + + if (max_neighbor_rank >= std::numeric_limits::max()) { + // Maximum rank neighbor is alreay in MIS + // Discard current vertex by setting its rank to + // std::numeric_limits::lowest() + ranks[v_offset] = std::numeric_limits::lowest(); + return true; + } + + if (rank_of_v >= max_neighbor_rank) { + // Include v and set its rank to std::numeric_limits::max() + ranks[v_offset] = std::numeric_limits::max(); + return true; + } + return false; + }); + + max_outgoing_ranks.resize(0, handle.get_stream()); + max_outgoing_ranks.shrink_to_fit(handle.get_stream()); + + d_sampled_vertices.resize(thrust::distance(d_sampled_vertices.begin(), last), + handle.get_stream()); + d_sampled_vertices.shrink_to_fit(handle.get_stream()); + + remaining_vertices.resize(non_candidate_vertices.size() + d_sampled_vertices.size(), + handle.get_stream()); + remaining_vertices.shrink_to_fit(handle.get_stream()); + + // merge non-candidate and remaining candidate vertices + thrust::merge(handle.get_thrust_policy(), + non_candidate_vertices.begin(), + non_candidate_vertices.end(), + d_sampled_vertices.begin(), + d_sampled_vertices.end(), + remaining_vertices.begin()); + + nr_remaining_vertices_to_check = remaining_vertices.size(); + if (multi_gpu) { + nr_remaining_vertices_to_check = host_scalar_allreduce(handle.get_comms(), + nr_remaining_vertices_to_check, + raft::comms::op_t::SUM, + handle.get_stream()); + } + + if (nr_remaining_vertices_to_check == 0) { break; } + } + + // Count number of vertices included in MIS + + vertex_t nr_vertices_included_in_mis = thrust::count_if( + handle.get_thrust_policy(), ranks.begin(), ranks.end(), [] __device__(auto v_rank) { + return v_rank >= std::numeric_limits::max(); + }); + + // Build MIS and return + rmm::device_uvector mis(nr_vertices_included_in_mis, handle.get_stream()); + thrust::copy_if( + handle.get_thrust_policy(), + vertex_begin, + vertex_end, + ranks.begin(), + mis.begin(), + [] __device__(auto v_rank) { return v_rank >= std::numeric_limits::max(); }); + + ranks.resize(0, handle.get_stream()); + ranks.shrink_to_fit(handle.get_stream()); + return mis; +} +} // namespace detail + +} // namespace cugraph diff --git a/cpp/src/community/mis.hpp b/cpp/src/community/detail/maximal_independent_moves.hpp similarity index 91% rename from cpp/src/community/mis.hpp rename to cpp/src/community/detail/maximal_independent_moves.hpp index 83c0d9775f9..b5588b11ef7 100644 --- a/cpp/src/community/mis.hpp +++ b/cpp/src/community/detail/maximal_independent_moves.hpp @@ -23,9 +23,13 @@ #include namespace cugraph { +namespace detail { + template -rmm::device_uvector maximal_independent_set( +rmm::device_uvector maximal_independent_moves( raft::handle_t const& handle, graph_view_t const& graph_view, raft::random::RngState& rng_state); + +} // namespace detail } // namespace cugraph diff --git a/cpp/src/community/detail/maximal_independent_moves_mg.cu b/cpp/src/community/detail/maximal_independent_moves_mg.cu new file mode 100644 index 00000000000..577253cdf58 --- /dev/null +++ b/cpp/src/community/detail/maximal_independent_moves_mg.cu @@ -0,0 +1,38 @@ +/* + * Copyright (c) 2023-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 "maximal_independent_moves.cuh" + +namespace cugraph { +namespace detail { + +template rmm::device_uvector maximal_independent_moves( + raft::handle_t const& handle, + graph_view_t const& decision_graph_view, + raft::random::RngState& rng_state); + +template rmm::device_uvector maximal_independent_moves( + raft::handle_t const& handle, + graph_view_t const& decision_graph_view, + raft::random::RngState& rng_state); + +template rmm::device_uvector maximal_independent_moves( + raft::handle_t const& handle, + graph_view_t const& decision_graph_view, + raft::random::RngState& rng_state); + +} // namespace detail + +} // namespace cugraph diff --git a/cpp/src/community/detail/maximal_independent_moves_sg.cu b/cpp/src/community/detail/maximal_independent_moves_sg.cu new file mode 100644 index 00000000000..18527c1ce48 --- /dev/null +++ b/cpp/src/community/detail/maximal_independent_moves_sg.cu @@ -0,0 +1,37 @@ +/* + * Copyright (c) 2023-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 "maximal_independent_moves.cuh" + +namespace cugraph { +namespace detail { + +template rmm::device_uvector maximal_independent_moves( + raft::handle_t const& handle, + graph_view_t const& decision_graph_view, + raft::random::RngState& rng_state); + +template rmm::device_uvector maximal_independent_moves( + raft::handle_t const& handle, + graph_view_t const& decision_graph_view, + raft::random::RngState& rng_state); + +template rmm::device_uvector maximal_independent_moves( + raft::handle_t const& handle, + graph_view_t const& decision_graph_view, + raft::random::RngState& rng_state); + +} // namespace detail +} // namespace cugraph diff --git a/cpp/src/community/detail/refine_impl.cuh b/cpp/src/community/detail/refine_impl.cuh index b767ce7d8bb..ef34ad90584 100644 --- a/cpp/src/community/detail/refine_impl.cuh +++ b/cpp/src/community/detail/refine_impl.cuh @@ -16,8 +16,8 @@ #pragma once #include "common_methods.hpp" -#include "community/mis.hpp" #include "detail/graph_partition_utils.cuh" +#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" @@ -660,8 +660,8 @@ refine_clustering( // Determine a set of moves using MIS of the decision_graph // - auto vertices_in_mis = - maximal_independent_set(handle, decision_graph_view, rng_state); + auto vertices_in_mis = maximal_independent_moves( + handle, decision_graph_view, rng_state); rmm::device_uvector numbering_indices((*renumber_map).size(), handle.get_stream()); detail::sequence_fill(handle.get_stream(), diff --git a/cpp/src/community/detail/mis_impl.cuh b/cpp/src/components/mis_impl.cuh similarity index 90% rename from cpp/src/community/detail/mis_impl.cuh rename to cpp/src/components/mis_impl.cuh index 217d64eb1c1..550edf9807a 100644 --- a/cpp/src/community/detail/mis_impl.cuh +++ b/cpp/src/components/mis_impl.cuh @@ -16,11 +16,12 @@ */ #pragma once -#include "community/mis.hpp" +#include "prims/fill_edge_property.cuh" #include "prims/fill_edge_src_dst_property.cuh" #include "prims/per_v_transform_reduce_incoming_outgoing_e.cuh" #include "prims/update_edge_src_dst_property.cuh" +#include #include #include #include @@ -60,36 +61,46 @@ rmm::device_uvector maximal_independent_set( thrust::make_counting_iterator(graph_view.local_vertex_partition_range_first()); auto vertex_end = thrust::make_counting_iterator(graph_view.local_vertex_partition_range_last()); - // Compute out-degree auto out_degrees = graph_view.compute_out_degrees(handle); + auto in_degrees = graph_view.compute_in_degrees(handle); - // Vertices with non-zero out-degree are possible candidates for MIS. + // Vertices with degree zero are always part of MIS remaining_vertices.resize( thrust::distance(remaining_vertices.begin(), thrust::copy_if(handle.get_thrust_policy(), vertex_begin, vertex_end, - out_degrees.begin(), + thrust::make_zip_iterator( + thrust::make_tuple(out_degrees.begin(), in_degrees.begin())), remaining_vertices.begin(), - [] __device__(auto deg) { return deg > 0; })), + [] __device__(auto out_deg_and_in_deg) { + return !((thrust::get<0>(out_deg_and_in_deg) == 0) && + (thrust::get<1>(out_deg_and_in_deg) == 0)); + })), handle.get_stream()); // Set ID of each vertex as its rank rmm::device_uvector ranks(local_vtx_partitoin_size, handle.get_stream()); thrust::copy(handle.get_thrust_policy(), vertex_begin, vertex_end, ranks.begin()); - // Set ranks of zero out-degree vetices to std::numeric_limits::lowest() - thrust::transform_if(handle.get_thrust_policy(), - out_degrees.begin(), - out_degrees.end(), - ranks.begin(), - cuda::proclaim_return_type( - [] __device__(auto) { return std::numeric_limits::lowest(); }), - [] __device__(auto deg) { return deg == 0; }); + // Set ranks of zero degree vetices to std::numeric_limits::max() + thrust::transform_if( + handle.get_thrust_policy(), + thrust::make_zip_iterator(thrust::make_tuple(out_degrees.begin(), in_degrees.begin())), + thrust::make_zip_iterator(thrust::make_tuple(out_degrees.end(), in_degrees.end())), + ranks.begin(), + cuda::proclaim_return_type( + [] __device__(auto) { return std::numeric_limits::max(); }), + [] __device__(auto in_out_degree) { + return (thrust::get<0>(in_out_degree) == 0) && (thrust::get<1>(in_out_degree) == 0); + }); out_degrees.resize(0, handle.get_stream()); out_degrees.shrink_to_fit(handle.get_stream()); + in_degrees.resize(0, handle.get_stream()); + in_degrees.shrink_to_fit(handle.get_stream()); + size_t loop_counter = 0; while (true) { loop_counter++; diff --git a/cpp/src/community/detail/mis_mg.cu b/cpp/src/components/mis_mg.cu similarity index 83% rename from cpp/src/community/detail/mis_mg.cu rename to cpp/src/components/mis_mg.cu index 0fc5eeb63c0..2418b38dd0b 100644 --- a/cpp/src/community/detail/mis_mg.cu +++ b/cpp/src/components/mis_mg.cu @@ -14,21 +14,21 @@ * limitations under the License. */ #include "mis_impl.cuh" - namespace cugraph { + template rmm::device_uvector maximal_independent_set( raft::handle_t const& handle, - graph_view_t const& decision_graph_view, + graph_view_t const& graph_view, raft::random::RngState& rng_state); template rmm::device_uvector maximal_independent_set( raft::handle_t const& handle, - graph_view_t const& decision_graph_view, + graph_view_t const& graph_view, raft::random::RngState& rng_state); template rmm::device_uvector maximal_independent_set( raft::handle_t const& handle, - graph_view_t const& decision_graph_view, + graph_view_t const& graph_view, raft::random::RngState& rng_state); } // namespace cugraph diff --git a/cpp/src/community/detail/mis_sg.cu b/cpp/src/components/mis_sg.cu similarity index 82% rename from cpp/src/community/detail/mis_sg.cu rename to cpp/src/components/mis_sg.cu index 8a20b31d157..fea4c4f2765 100644 --- a/cpp/src/community/detail/mis_sg.cu +++ b/cpp/src/components/mis_sg.cu @@ -16,19 +16,20 @@ #include "mis_impl.cuh" namespace cugraph { + template rmm::device_uvector maximal_independent_set( raft::handle_t const& handle, - graph_view_t const& decision_graph_view, + graph_view_t const& graph_view, raft::random::RngState& rng_state); template rmm::device_uvector maximal_independent_set( raft::handle_t const& handle, - graph_view_t const& decision_graph_view, + graph_view_t const& graph_view, raft::random::RngState& rng_state); template rmm::device_uvector maximal_independent_set( raft::handle_t const& handle, - graph_view_t const& decision_graph_view, + graph_view_t const& graph_view, raft::random::RngState& rng_state); } // namespace cugraph diff --git a/cpp/src/components/vertex_coloring_impl.cuh b/cpp/src/components/vertex_coloring_impl.cuh new file mode 100644 index 00000000000..ce445ab3809 --- /dev/null +++ b/cpp/src/components/vertex_coloring_impl.cuh @@ -0,0 +1,150 @@ +/* + * 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/transform_e.cuh" +#include "prims/update_edge_src_dst_property.cuh" + +#include + +#include +#include + +namespace cugraph { + +namespace detail { + +template +rmm::device_uvector vertex_coloring( + raft::handle_t const& handle, + cugraph::graph_view_t const& graph_view, + raft::random::RngState& rng_state) +{ + using graph_view_t = cugraph::graph_view_t; + graph_view_t current_graph_view(graph_view); + + // 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); + + cugraph::transform_e( + handle, + current_graph_view, + edge_src_dummy_property_t{}.view(), + 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); // mask out self-loop + }, + edge_masks_even.mutable_view()); + + current_graph_view.attach_edge_mask(edge_masks_even.view()); + + // device vector to store colors of vertices + rmm::device_uvector colors = rmm::device_uvector( + current_graph_view.local_vertex_partition_range_size(), handle.get_stream()); + thrust::fill( + handle.get_thrust_policy(), colors.begin(), colors.end(), std::numeric_limits::max()); + + vertex_t color_id = 0; + while (true) { + auto mis = cugraph::maximal_independent_set( + handle, current_graph_view, rng_state); + + using flag_t = uint8_t; + rmm::device_uvector is_vertex_in_mis = rmm::device_uvector( + current_graph_view.local_vertex_partition_range_size(), handle.get_stream()); + thrust::fill(handle.get_thrust_policy(), is_vertex_in_mis.begin(), is_vertex_in_mis.end(), 0); + + thrust::for_each( + handle.get_thrust_policy(), + mis.begin(), + mis.end(), + [color_id, + colors = colors.data(), + is_vertex_in_mis = is_vertex_in_mis.data(), + v_first = current_graph_view.local_vertex_partition_range_first()] __device__(vertex_t v) { + auto v_offset = v - v_first; + is_vertex_in_mis[v_offset] = flag_t{1}; + vertex_t initial_color_id = colors[v_offset]; + colors[v_offset] = (color_id < initial_color_id) ? color_id : initial_color_id; + }); + + if (current_graph_view.compute_number_of_edges(handle) == 0) { break; } + + cugraph::edge_src_property_t src_mis_flags(handle, current_graph_view); + cugraph::edge_dst_property_t dst_mis_flags(handle, current_graph_view); + + cugraph::update_edge_src_property( + handle, current_graph_view, is_vertex_in_mis.begin(), src_mis_flags); + + cugraph::update_edge_dst_property( + handle, current_graph_view, is_vertex_in_mis.begin(), dst_mis_flags); + + if (color_id % 2 == 0) { + cugraph::transform_e( + handle, + current_graph_view, + src_mis_flags.view(), + dst_mis_flags.view(), + cugraph::edge_dummy_property_t{}.view(), + [color_id] __device__( + auto src, auto dst, auto is_src_in_mis, auto is_dst_in_mis, thrust::nullopt_t) { + return !((is_src_in_mis == uint8_t{true}) || (is_dst_in_mis == uint8_t{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 { + cugraph::transform_e( + handle, + current_graph_view, + src_mis_flags.view(), + dst_mis_flags.view(), + cugraph::edge_dummy_property_t{}.view(), + [color_id] __device__( + auto src, auto dst, auto is_src_in_mis, auto is_dst_in_mis, thrust::nullopt_t) { + return !((is_src_in_mis == uint8_t{true}) || (is_dst_in_mis == uint8_t{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()); + } + + color_id++; + } + return colors; +} +} // namespace detail + +template +rmm::device_uvector vertex_coloring( + raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::random::RngState& rng_state) +{ + return detail::vertex_coloring(handle, graph_view, rng_state); +} + +} // namespace cugraph diff --git a/cpp/src/components/vertex_coloring_mg.cu b/cpp/src/components/vertex_coloring_mg.cu new file mode 100644 index 00000000000..8f87e8bd534 --- /dev/null +++ b/cpp/src/components/vertex_coloring_mg.cu @@ -0,0 +1,35 @@ +/* + * 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 "vertex_coloring_impl.cuh" + +namespace cugraph { + +template rmm::device_uvector vertex_coloring( + raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::random::RngState& rng_state); + +template rmm::device_uvector vertex_coloring( + raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::random::RngState& rng_state); + +template rmm::device_uvector vertex_coloring( + raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::random::RngState& rng_state); + +} // namespace cugraph diff --git a/cpp/src/components/vertex_coloring_sg.cu b/cpp/src/components/vertex_coloring_sg.cu new file mode 100644 index 00000000000..427bc0b2c81 --- /dev/null +++ b/cpp/src/components/vertex_coloring_sg.cu @@ -0,0 +1,35 @@ +/* + * 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 "vertex_coloring_impl.cuh" + +namespace cugraph { + +template rmm::device_uvector vertex_coloring( + raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::random::RngState& rng_state); + +template rmm::device_uvector vertex_coloring( + raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::random::RngState& rng_state); + +template rmm::device_uvector vertex_coloring( + raft::handle_t const& handle, + graph_view_t const& graph_view, + raft::random::RngState& rng_state); + +} // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 46a895536ef..af0dffcbf65 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -372,6 +372,14 @@ ConfigureTest(EDGE_BETWEENNESS_CENTRALITY_TEST centrality/edge_betweenness_centr # - WEAKLY CONNECTED COMPONENTS tests ------------------------------------------------------------- ConfigureTest(WEAKLY_CONNECTED_COMPONENTS_TEST components/weakly_connected_components_test.cpp) +############################################################################################### +# - MIS tests ------------------------------------------------------------------------------ +ConfigureTest(MIS_TEST components/mis_test.cu) + +############################################################################################### +# - VERTEX COLORING tests ------------------------------------------------------------------- +ConfigureTest(VERTEX_COLORING_TEST components/vertex_coloring_test.cu) + ################################################################################################### # - SIMILARITY tests ------------------------------------------------------------------------------ ConfigureTest(SIMILARITY_TEST link_prediction/similarity_test.cu) @@ -535,10 +543,6 @@ if(BUILD_CUGRAPH_MG_TESTS) # - MG ECG tests -------------------------------------------------------------------------- ConfigureTestMG(MG_ECG_TEST community/mg_ecg_test.cpp) - ############################################################################################### - # - MG MIS tests ------------------------------------------------------------------------------ - ConfigureTestMG(MG_MIS_TEST community/mg_mis_test.cu) - ############################################################################################### # - MG SELECT RANDOM VERTICES tests ----------------------------------------------------------- ConfigureTestMG(MG_SELECT_RANDOM_VERTICES structure/mg_select_random_vertices_test.cpp) @@ -552,6 +556,14 @@ if(BUILD_CUGRAPH_MG_TESTS) ConfigureTestMG(MG_WEAKLY_CONNECTED_COMPONENTS_TEST components/mg_weakly_connected_components_test.cpp) + ############################################################################################### + # - MG MIS tests ------------------------------------------------------------------------------ + ConfigureTestMG(MG_MIS_TEST components/mg_mis_test.cu) + + ############################################################################################### + # - MG VERTEX COLORING tests ------------------------------------------------------------------- + ConfigureTestMG(MG_VERTEX_COLORING_TEST components/mg_vertex_coloring_test.cu) + ############################################################################################### # - MG Core Number tests ---------------------------------------------------------------------- ConfigureTestMG(MG_CORE_NUMBER_TEST cores/mg_core_number_test.cpp) diff --git a/cpp/tests/community/mg_mis_test.cu b/cpp/tests/components/mg_mis_test.cu similarity index 95% rename from cpp/tests/community/mg_mis_test.cu rename to cpp/tests/components/mg_mis_test.cu index 1240cf812f9..04c346b0f00 100644 --- a/cpp/tests/community/mg_mis_test.cu +++ b/cpp/tests/components/mg_mis_test.cu @@ -245,18 +245,20 @@ TEST_P(Tests_MGMaximalIndependentSet_Rmat, CheckInt64Int64FloatFloat) override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); } +bool constexpr check_correctness = false; INSTANTIATE_TEST_SUITE_P( file_test, Tests_MGMaximalIndependentSet_File, - ::testing::Combine(::testing::Values(MaximalIndependentSet_Usecase{false}, - MaximalIndependentSet_Usecase{false}), + ::testing::Combine(::testing::Values(MaximalIndependentSet_Usecase{check_correctness}, + MaximalIndependentSet_Usecase{check_correctness}), ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); -INSTANTIATE_TEST_SUITE_P(rmat_small_test, - Tests_MGMaximalIndependentSet_Rmat, - ::testing::Combine(::testing::Values(MaximalIndependentSet_Usecase{false}), - ::testing::Values(cugraph::test::Rmat_Usecase( - 3, 4, 0.57, 0.19, 0.19, 0, true, false)))); +INSTANTIATE_TEST_SUITE_P( + rmat_small_test, + Tests_MGMaximalIndependentSet_Rmat, + ::testing::Combine( + ::testing::Values(MaximalIndependentSet_Usecase{check_correctness}), + ::testing::Values(cugraph::test::Rmat_Usecase(3, 4, 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 @@ -266,7 +268,8 @@ INSTANTIATE_TEST_SUITE_P( factor (to avoid running same benchmarks more than once) */ Tests_MGMaximalIndependentSet_Rmat, ::testing::Combine( - ::testing::Values(MaximalIndependentSet_Usecase{false}, MaximalIndependentSet_Usecase{false}), + ::testing::Values(MaximalIndependentSet_Usecase{check_correctness}, + MaximalIndependentSet_Usecase{check_correctness}), ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/components/mg_vertex_coloring_test.cu b/cpp/tests/components/mg_vertex_coloring_test.cu new file mode 100644 index 00000000000..ce4dadaa786 --- /dev/null +++ b/cpp/tests/components/mg_vertex_coloring_test.cu @@ -0,0 +1,275 @@ +/* + * 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 "prims/fill_edge_src_dst_property.cuh" +#include "prims/per_v_transform_reduce_incoming_outgoing_e.cuh" +#include "prims/property_generator.cuh" +#include "prims/reduce_op.cuh" +#include "prims/transform_reduce_e.cuh" +#include "prims/update_edge_src_dst_property.cuh" +#include "utilities/base_fixture.hpp" +#include "utilities/test_graphs.hpp" +#include "utilities/test_utilities.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +struct GraphColoring_UseCase { + bool check_correctness{true}; +}; + +template +class Tests_MGGraphColoring + : public ::testing::TestWithParam> { + public: + Tests_MGGraphColoring() {} + + 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 [coloring_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; + + auto [mg_graph, mg_edge_weights, mg_renumber_map] = + cugraph::test::construct_graph( + *handle_, input_usecase, false, true); + + 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; + + raft::random::RngState rng_state(multi_gpu ? handle_->get_comms().get_rank() : 0); + auto d_colors = + cugraph::vertex_coloring(*handle_, mg_graph_view, rng_state); + + // Test Graph Coloring + + if (coloring_usecase.check_correctness) { + std::vector h_colors(d_colors.size()); + raft::update_host(h_colors.data(), d_colors.data(), d_colors.size(), handle_->get_stream()); + + std::for_each(h_colors.begin(), + h_colors.end(), + [num_vertices = mg_graph_view.number_of_vertices()](vertex_t color_id) { + ASSERT_TRUE(color_id <= num_vertices); + }); + + using GraphViewType = cugraph::graph_view_t; + cugraph::edge_src_property_t src_color_cache(*handle_); + cugraph::edge_dst_property_t dst_color_cache(*handle_); + + if constexpr (multi_gpu) { + src_color_cache = + cugraph::edge_src_property_t(*handle_, mg_graph_view); + dst_color_cache = + cugraph::edge_dst_property_t(*handle_, mg_graph_view); + update_edge_src_property(*handle_, mg_graph_view, d_colors.begin(), src_color_cache); + update_edge_dst_property(*handle_, mg_graph_view, d_colors.begin(), dst_color_cache); + } + + rmm::device_uvector d_color_conflicts( + mg_graph_view.local_vertex_partition_range_size(), handle_->get_stream()); + + per_v_transform_reduce_outgoing_e( + *handle_, + mg_graph_view, + multi_gpu + ? src_color_cache.view() + : cugraph::detail::edge_major_property_view_t(d_colors.data()), + multi_gpu ? dst_color_cache.view() + : cugraph::detail::edge_minor_property_view_t( + d_colors.data(), vertex_t{0}), + cugraph::edge_dummy_property_t{}.view(), + [] __device__(auto src, auto dst, auto src_color, auto dst_color, thrust::nullopt_t) { + if ((src != dst) && (src_color == dst_color)) { + return uint8_t{1}; + } else { + return uint8_t{0}; + } + }, + uint8_t{0}, + cugraph::reduce_op::maximum{}, + d_color_conflicts.begin()); + + std::vector h_color_conflicts(d_color_conflicts.size()); + raft::update_host(h_color_conflicts.data(), + d_color_conflicts.data(), + d_color_conflicts.size(), + handle_->get_stream()); + + std::vector h_vertices_in_this_proces((*mg_renumber_map).size()); + + raft::update_host(h_vertices_in_this_proces.data(), + (*mg_renumber_map).data(), + (*mg_renumber_map).size(), + handle_->get_stream()); + handle_->sync_stream(); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + size_t nr_conflicts = cugraph::transform_reduce_e( + *handle_, + mg_graph_view, + multi_gpu ? src_color_cache.view() + : cugraph::detail::edge_major_property_view_t( + d_colors.begin()), + multi_gpu ? dst_color_cache.view() + : cugraph::detail::edge_minor_property_view_t( + d_colors.begin(), vertex_t{0}), + cugraph::edge_dummy_property_t{}.view(), + [renumber_map = (*mg_renumber_map).data()] __device__( + auto src, auto dst, auto src_color, auto dst_color, thrust::nullopt_t) { + if ((src != dst) && (src_color == dst_color)) { + return vertex_t{1}; + } else { + return vertex_t{0}; + } + }, + vertex_t{0}); + + ASSERT_TRUE(nr_conflicts == edge_t{0}) + << "adjacent vertices can't have same color." << std::endl; + + { + thrust::for_each( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple( + h_colors.begin(), h_vertices_in_this_proces.begin(), h_color_conflicts.begin())), + thrust::make_zip_iterator(thrust::make_tuple( + h_colors.end(), h_vertices_in_this_proces.end(), h_color_conflicts.end())), + [](auto color_vetex_and_conflict_flag) { + auto color = thrust::get<0>(color_vetex_and_conflict_flag); + auto v = thrust::get<1>(color_vetex_and_conflict_flag); + auto conflict_flag = thrust::get<2>(color_vetex_and_conflict_flag); + ASSERT_TRUE(conflict_flag == 0) + << v << " got same color as one of its neighbor" << std::endl; + }); + } + } + } + + private: + static std::unique_ptr handle_; +}; + +template +std::unique_ptr Tests_MGGraphColoring::handle_ = nullptr; + +using Tests_MGGraphColoring_File = Tests_MGGraphColoring; +using Tests_MGGraphColoring_Rmat = Tests_MGGraphColoring; + +TEST_P(Tests_MGGraphColoring_File, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGGraphColoring_File, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGGraphColoring_File, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGGraphColoring_Rmat, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGGraphColoring_Rmat, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_MGGraphColoring_Rmat, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +bool constexpr check_correctness = false; + +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_MGGraphColoring_File, + ::testing::Combine(::testing::Values(GraphColoring_UseCase{check_correctness}, + GraphColoring_UseCase{check_correctness}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + rmat_small_test, + Tests_MGGraphColoring_Rmat, + ::testing::Combine( + ::testing::Values(GraphColoring_UseCase{check_correctness}), + ::testing::Values(cugraph::test::Rmat_Usecase(3, 4, 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_MGGraphColoring_Rmat, + ::testing::Combine( + ::testing::Values(GraphColoring_UseCase{check_correctness}, + GraphColoring_UseCase{check_correctness}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); + +CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/components/mis_test.cu b/cpp/tests/components/mis_test.cu new file mode 100644 index 00000000000..f3bdd3d0e8b --- /dev/null +++ b/cpp/tests/components/mis_test.cu @@ -0,0 +1,240 @@ +/* + * 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 "prims/per_v_transform_reduce_incoming_outgoing_e.cuh" +#include "prims/reduce_op.cuh" +#include "utilities/base_fixture.hpp" +#include "utilities/test_graphs.hpp" +#include "utilities/test_utilities.hpp" + +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +struct MaximalIndependentSet_Usecase { + bool check_correctness{true}; +}; + +template +class Tests_SGMaximalIndependentSet + : public ::testing::TestWithParam> { + public: + Tests_SGMaximalIndependentSet() {} + + static void SetUpTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [mis_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; + + auto [sg_graph, sg_edge_weights, sg_renumber_map] = + cugraph::test::construct_graph( + handle, input_usecase, false, true); + + 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; + + raft::random::RngState rng_state(0); + auto d_mis = cugraph::maximal_independent_set( + handle, sg_graph_view, rng_state); + + // Test MIS + if (mis_usecase.check_correctness) { + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + std::vector h_mis(d_mis.size()); + raft::update_host(h_mis.data(), d_mis.data(), d_mis.size(), handle.get_stream()); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + auto vertex_first = sg_graph_view.local_vertex_partition_range_first(); + auto vertex_last = sg_graph_view.local_vertex_partition_range_last(); + + std::for_each(h_mis.begin(), h_mis.end(), [vertex_first, vertex_last](vertex_t v) { + ASSERT_TRUE((v >= vertex_first) && (v < vertex_last)); + }); + + // If a vertex is included in MIS, then none of its neighbor should be + + vertex_t local_vtx_partitoin_size = sg_graph_view.local_vertex_partition_range_size(); + rmm::device_uvector d_total_outgoing_nbrs_included_mis(local_vtx_partitoin_size, + handle.get_stream()); + + rmm::device_uvector inclusiong_flags(local_vtx_partitoin_size, handle.get_stream()); + + thrust::uninitialized_fill( + handle.get_thrust_policy(), inclusiong_flags.begin(), inclusiong_flags.end(), vertex_t{0}); + + thrust::for_each( + handle.get_thrust_policy(), + d_mis.begin(), + d_mis.end(), + [inclusiong_flags = + raft::device_span(inclusiong_flags.data(), inclusiong_flags.size()), + v_first = sg_graph_view.local_vertex_partition_range_first()] __device__(auto v) { + auto v_offset = v - v_first; + inclusiong_flags[v_offset] = vertex_t{1}; + }); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + per_v_transform_reduce_outgoing_e( + handle, + sg_graph_view, + cugraph::detail::edge_major_property_view_t( + inclusiong_flags.data()), + cugraph::detail::edge_minor_property_view_t( + inclusiong_flags.data(), vertex_t{0}), + 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; + }, + vertex_t{0}, + cugraph::reduce_op::plus{}, + d_total_outgoing_nbrs_included_mis.begin()); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + std::vector h_total_outgoing_nbrs_included_mis( + d_total_outgoing_nbrs_included_mis.size()); + raft::update_host(h_total_outgoing_nbrs_included_mis.data(), + d_total_outgoing_nbrs_included_mis.data(), + d_total_outgoing_nbrs_included_mis.size(), + handle.get_stream()); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + { + auto vertex_first = sg_graph_view.local_vertex_partition_range_first(); + auto vertex_last = sg_graph_view.local_vertex_partition_range_last(); + + std::for_each(h_mis.begin(), + h_mis.end(), + [vertex_first, vertex_last, &h_total_outgoing_nbrs_included_mis](vertex_t v) { + ASSERT_TRUE((v >= vertex_first) && (v < vertex_last)) + << v << " is not within vertex parition range" << std::endl; + + ASSERT_TRUE(h_total_outgoing_nbrs_included_mis[v - vertex_first] == 0) + << v << "'s neighbor is included in MIS" << std::endl; + }); + } + } + } +}; + +using Tests_SGMaximalIndependentSet_File = + Tests_SGMaximalIndependentSet; +using Tests_SGMaximalIndependentSet_Rmat = + Tests_SGMaximalIndependentSet; + +TEST_P(Tests_SGMaximalIndependentSet_File, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGMaximalIndependentSet_File, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGMaximalIndependentSet_File, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGMaximalIndependentSet_Rmat, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGMaximalIndependentSet_Rmat, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGMaximalIndependentSet_Rmat, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +bool constexpr check_correctness = false; +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_SGMaximalIndependentSet_File, + ::testing::Combine(::testing::Values(MaximalIndependentSet_Usecase{check_correctness}, + MaximalIndependentSet_Usecase{check_correctness}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + rmat_small_test, + Tests_SGMaximalIndependentSet_Rmat, + ::testing::Combine( + ::testing::Values(MaximalIndependentSet_Usecase{check_correctness}), + ::testing::Values(cugraph::test::Rmat_Usecase(3, 4, 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_SGMaximalIndependentSet_Rmat, + ::testing::Combine( + ::testing::Values(MaximalIndependentSet_Usecase{check_correctness}, + MaximalIndependentSet_Usecase{check_correctness}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); + +CUGRAPH_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/components/vertex_coloring_test.cu b/cpp/tests/components/vertex_coloring_test.cu new file mode 100644 index 00000000000..27a0c5013bd --- /dev/null +++ b/cpp/tests/components/vertex_coloring_test.cu @@ -0,0 +1,244 @@ +/* + * 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 "prims/per_v_transform_reduce_incoming_outgoing_e.cuh" +#include "prims/reduce_op.cuh" +#include "prims/transform_reduce_e.cuh" +#include "utilities/base_fixture.hpp" +#include "utilities/test_graphs.hpp" +#include "utilities/test_utilities.hpp" + +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +struct GraphColoring_UseCase { + bool check_correctness{true}; +}; + +template +class Tests_SGGraphColoring + : public ::testing::TestWithParam> { + public: + Tests_SGGraphColoring() {} + + static void SetUpTestCase() {} + static void TearDownTestCase() {} + + virtual void SetUp() {} + virtual void TearDown() {} + + template + void run_current_test(std::tuple const& param) + { + auto [coloring_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; + + auto [sg_graph, sg_edge_weights, sg_renumber_map] = + cugraph::test::construct_graph( + handle, input_usecase, false, true); + + 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; + + raft::random::RngState rng_state(0); + auto d_colors = + cugraph::vertex_coloring(handle, sg_graph_view, rng_state); + + // Test Graph Coloring + + if (coloring_usecase.check_correctness) { + std::vector h_colors(d_colors.size()); + raft::update_host(h_colors.data(), d_colors.data(), d_colors.size(), handle.get_stream()); + + std::for_each(h_colors.begin(), + h_colors.end(), + [num_vertices = sg_graph_view.number_of_vertices()](vertex_t color_id) { + ASSERT_TRUE(color_id <= num_vertices); + }); + + rmm::device_uvector d_color_conflict_flags( + sg_graph_view.local_vertex_partition_range_size(), handle.get_stream()); + + per_v_transform_reduce_outgoing_e( + handle, + sg_graph_view, + cugraph::detail::edge_major_property_view_t(d_colors.data()), + cugraph::detail::edge_minor_property_view_t(d_colors.data(), + vertex_t{0}), + cugraph::edge_dummy_property_t{}.view(), + [] __device__(auto src, auto dst, auto src_color, auto dst_color, thrust::nullopt_t) { + if ((src != dst) && (src_color == dst_color)) { + return uint8_t{1}; + } else { + return uint8_t{0}; + } + }, + uint8_t{0}, + cugraph::reduce_op::maximum{}, + d_color_conflict_flags.begin()); + + std::vector h_color_conflict_flags(d_color_conflict_flags.size()); + raft::update_host(h_color_conflict_flags.data(), + d_color_conflict_flags.data(), + d_color_conflict_flags.size(), + handle.get_stream()); + + std::vector h_vertices_in_this_proces((*sg_renumber_map).size()); + + raft::update_host(h_vertices_in_this_proces.data(), + (*sg_renumber_map).data(), + (*sg_renumber_map).size(), + handle.get_stream()); + handle.sync_stream(); + + RAFT_CUDA_TRY(cudaDeviceSynchronize()); + + edge_t nr_conflicts = cugraph::transform_reduce_e( + handle, + sg_graph_view, + cugraph::detail::edge_major_property_view_t(d_colors.begin()), + cugraph::detail::edge_minor_property_view_t(d_colors.begin(), + vertex_t{0}), + cugraph::edge_dummy_property_t{}.view(), + [renumber_map = (*sg_renumber_map).data()] __device__( + auto src, auto dst, auto src_color, auto dst_color, thrust::nullopt_t) { + if ((src != dst) && (src_color == dst_color)) { + return vertex_t{1}; + } else { + return vertex_t{0}; + } + }, + vertex_t{0}); + + ASSERT_TRUE(nr_conflicts == edge_t{0}) + << "adjacent vertices can't have same color." << std::endl; + + if (nr_conflicts >= 0) { + thrust::for_each( + thrust::host, + thrust::make_zip_iterator(thrust::make_tuple( + h_colors.begin(), h_vertices_in_this_proces.begin(), h_color_conflict_flags.begin())), + thrust::make_zip_iterator(thrust::make_tuple( + h_colors.end(), h_vertices_in_this_proces.end(), h_color_conflict_flags.end())), + [](auto color_vetex_and_conflict_flag) { + auto color = thrust::get<0>(color_vetex_and_conflict_flag); + auto v = thrust::get<1>(color_vetex_and_conflict_flag); + auto conflict_flag = thrust::get<2>(color_vetex_and_conflict_flag); + ASSERT_TRUE(conflict_flag == 0) + << v << " got same color as one of its neighbor" << std::endl; + }); + } + } + } +}; + +using Tests_SGGraphColoring_File = Tests_SGGraphColoring; +using Tests_SGGraphColoring_Rmat = Tests_SGGraphColoring; + +TEST_P(Tests_SGGraphColoring_File, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGGraphColoring_File, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGGraphColoring_File, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_File_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGGraphColoring_Rmat, CheckInt32Int32FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGGraphColoring_Rmat, CheckInt32Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +TEST_P(Tests_SGGraphColoring_Rmat, CheckInt64Int64FloatFloat) +{ + run_current_test( + override_Rmat_Usecase_with_cmd_line_arguments(GetParam())); +} + +bool constexpr check_correctness = false; + +INSTANTIATE_TEST_SUITE_P( + file_test, + Tests_SGGraphColoring_File, + ::testing::Combine(::testing::Values(GraphColoring_UseCase{check_correctness}, + GraphColoring_UseCase{check_correctness}), + ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx")))); + +INSTANTIATE_TEST_SUITE_P( + rmat_small_test, + Tests_SGGraphColoring_Rmat, + ::testing::Combine( + ::testing::Values(GraphColoring_UseCase{check_correctness}), + ::testing::Values(cugraph::test::Rmat_Usecase(3, 4, 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_SGGraphColoring_Rmat, + ::testing::Combine( + ::testing::Values(GraphColoring_UseCase{check_correctness}, + GraphColoring_UseCase{check_correctness}), + ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); + +CUGRAPH_TEST_PROGRAM_MAIN()