MG Leiden and MG MIS

cpp/include/cugraph/algorithms.hpp

@@ -760,7 +760,8 @@ std::pair<std::unique_ptr<Dendrogram<vertex_t>>, weight_t> leiden(
  graph_view_t<vertex_t, edge_t, false, multi_gpu> const& graph_view,
  std::optional<edge_property_view_t<edge_t, weight_t const*>> edge_weight_view,
  size_t max_level = 100,
- weight_t resolution = weight_t{1});
+ weight_t resolution = weight_t{1},
+ weight_t theta = weight_t{1});
 
 /**
  * @brief Leiden implementation
@@ -804,7 +805,8 @@ std::pair<size_t, weight_t> leiden(
  std::optional<edge_property_view_t<edge_t, weight_t const*>> edge_weight_view,
  vertex_t* clustering, // FIXME: Use (device_)span instead
  size_t max_level = 100,
- weight_t resolution = weight_t{1});
+ weight_t resolution = weight_t{1},
+ weight_t theta = weight_t{1});
 
 /**
  * @brief Computes the ecg clustering of the given graph.
@@ -2033,6 +2035,13 @@ std::tuple<rmm::device_uvector<size_t>, rmm::device_uvector<vertex_t>> k_hop_nbr
  size_t k,
  bool do_expensive_check = false);
 
+// FIXME: Write doc string
+template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
+rmm::device_uvector<vertex_t> compute_mis(
+ raft::handle_t const& handle,
+ graph_view_t<vertex_t, edge_t, false, multi_gpu> const& graph_view,
+ std::optional<edge_property_view_t<edge_t, weight_t const*>> edge_weight_view);
+
 } // namespace cugraph
 
 /**

cpp/src/community/detail/mis_impl.cuh

@@ -16,38 +16,28 @@
  */
 #pragma once
 
+#include <community/mis.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 <community/detail/mis.hpp>
 #include <cugraph/edge_property.hpp>
 #include <cugraph/edge_src_dst_property.hpp>
 #include <cugraph/graph_functions.hpp>
 #include <cugraph/graph_view.hpp>
 #include <cugraph/utilities/host_scalar_comm.hpp>
 
-#include <raft/util/cudart_utils.hpp>
-#include <raft/util/integer_utils.hpp>
-#include <rmm/exec_policy.hpp>
-
 #include <thrust/count.h>
 #include <thrust/distance.h>
-#include <thrust/execution_policy.h>
-#include <thrust/fill.h>
-#include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
+#include <thrust/merge.h>
 #include <thrust/optional.h>
 #include <thrust/remove.h>
-#include <thrust/sequence.h>
-#include <thrust/shuffle.h>
+#include <thrust/set_operations.h>
 #include <thrust/transform.h>
 #include <thrust/transform_reduce.h>
 
 #include <cmath>
-#include <numeric>
-#include <type_traits>
-#include <utility>
 
 namespace cugraph {
 
@@ -102,10 +92,7 @@ rmm::device_uvector<vertex_t> compute_mis(
  out_degrees.resize(0, handle.get_stream());
  out_degrees.shrink_to_fit(handle.get_stream());
 
- thrust::default_random_engine g;
- size_t seed = 0;
- if constexpr (multi_gpu) { seed = handle.get_comms().get_rank(); }
- g.seed(seed);
+ raft::random::RngState rng_state(multi_gpu ? handle.get_comms().get_rank() : 0);
 
  size_t loop_counter = 0;
  while (true) {
@@ -117,22 +104,47 @@ rmm::device_uvector<vertex_t> compute_mis(
  thrust::copy(handle.get_thrust_policy(), ranks.begin(), ranks.end(), temporary_ranks.begin());
 
  // Select a random set of candidate vertices
- // FIXME: use common utility function to select a subset of remaining vertices
- // and for MG extension, select from disributed array remaining vertices
- thrust::shuffle(
- handle.get_thrust_policy(), remaining_vertices.begin(), remaining_vertices.end(), g);
 
- vertex_t nr_candidates =
- (remaining_vertices.size() < 1024)
- ? remaining_vertices.size()
- : std::min(static_cast<vertex_t>((0.50 + 0.25 * loop_counter) * remaining_vertices.size()),
- static_cast<vertex_t>(remaining_vertices.size()));
+ 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<vertex_t>((0.50 + 0.25 * loop_counter) *
+ nr_remaining_vertices_to_check),
+ nr_remaining_vertices_to_check);
+
+ // FIXME: (nr_remaining_vertices_to_check < 1024)? avoid calling select_random_vertices
+ auto d_sampled_vertices =
+ cugraph::select_random_vertices(handle,
+ graph_view,
+ std::make_optional(raft::device_span<vertex_t const>{
+ remaining_vertices.data(), remaining_vertices.size()}),
+ rng_state,
+ nr_candidates,
+ false,
+ true);
+
+ rmm::device_uvector<vertex_t> 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<vertex_t>::lowest()
  thrust::for_each(
  handle.get_thrust_policy(),
- remaining_vertices.begin(),
- remaining_vertices.end() - nr_candidates,
+ non_candidate_vertices.begin(),
+ non_candidate_vertices.end(),
  [temporary_ranks =
  raft::device_span<vertex_t>(temporary_ranks.data(), temporary_ranks.size()),
  v_first = graph_view.local_vertex_partition_range_first()] __device__(auto v) {
@@ -224,8 +236,8 @@ rmm::device_uvector<vertex_t> compute_mis(
  //
  auto last = thrust::remove_if(
  handle.get_thrust_policy(),
- remaining_vertices.end() - nr_candidates,
- remaining_vertices.end(),
+ d_sampled_vertices.begin(),
+ d_sampled_vertices.end(),
  [max_rank_neighbor_first = max_outgoing_ranks.begin(),
  ranks = raft::device_span<vertex_t>(ranks.data(), ranks.size()),
  v_first = graph_view.local_vertex_partition_range_first()] __device__(auto v) {
@@ -252,11 +264,23 @@ rmm::device_uvector<vertex_t> compute_mis(
  max_outgoing_ranks.resize(0, handle.get_stream());
  max_outgoing_ranks.shrink_to_fit(handle.get_stream());
 
- remaining_vertices.resize(thrust::distance(remaining_vertices.begin(), last),
+ 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());
 
- vertex_t nr_remaining_vertices_to_check = remaining_vertices.size();
+ // 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,
@@ -289,4 +313,14 @@ rmm::device_uvector<vertex_t> compute_mis(
  return mis;
 }
 } // namespace detail
+
+template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
+rmm::device_uvector<vertex_t> compute_mis(
+ raft::handle_t const& handle,
+ graph_view_t<vertex_t, edge_t, false, multi_gpu> const& graph_view,
+ std::optional<edge_property_view_t<edge_t, weight_t const*>> edge_weight_view)
+{
+ return detail::compute_mis(handle, graph_view, edge_weight_view);
+}
+
 } // namespace cugraph

cpp/src/community/detail/mis_mg.cu

@@ -16,7 +16,6 @@
 #include <community/detail/mis_impl.cuh>
 
 namespace cugraph {
-namespace detail {
 template rmm::device_uvector<int32_t> compute_mis(
  raft::handle_t const& handle,
  graph_view_t<int32_t, int32_t, false, true> const& decision_graph_view,
@@ -47,5 +46,4 @@ template rmm::device_uvector<int64_t> compute_mis(
  graph_view_t<int64_t, int64_t, false, true> const& decision_graph_view,
  std::optional<edge_property_view_t<int64_t, double const*>> edge_weight_view);
 
-} // namespace detail
 } // namespace cugraph

cpp/src/community/detail/mis_sg.cu

@@ -16,7 +16,6 @@
 #include <community/detail/mis_impl.cuh>
 
 namespace cugraph {
-namespace detail {
 template rmm::device_uvector<int32_t> compute_mis(
  raft::handle_t const& handle,
  graph_view_t<int32_t, int32_t, false, false> const& decision_graph_view,
@@ -47,5 +46,4 @@ template rmm::device_uvector<int64_t> compute_mis(
  graph_view_t<int64_t, int64_t, false, false> const& decision_graph_view,
  std::optional<edge_property_view_t<int64_t, double const*>> edge_weight_view);
 
-} // namespace detail
 } // namespace cugraph

cpp/src/community/detail/refine.hpp

@@ -36,6 +36,7 @@ refine_clustering(
  edge_weight_view,
  weight_t total_edge_weight,
  weight_t resolution,
+ weight_t theta,
  rmm::device_uvector<weight_t> const& vertex_weights_v,
  rmm::device_uvector<typename graph_view_t::vertex_type>&& cluster_keys_v,
  rmm::device_uvector<weight_t>&& cluster_weights_v,