Update Louvain to use new graph primitives and pattern accelerators (#…

…1423) Implement the `update_by_delta_modularity` method using the new graph primitives and pattern accelerators. This eliminates all of the custom MNMG implementation originally created for MNMG Louvain a few releases ago and replaces it with the new pattern accelerator and graph primitives that have been added in the last couple of releases. This depends on the following PRs and should not be merged until after them: * #1394 * #1399 closes #1220 Authors: - Chuck Hastings (@ChuckHastings) Approvers: - Andrei Schaffer (@aschaffer) - Seunghwa Kang (@seunghwak) - Rick Ratzel (@rlratzel) - Alex Fender (@afender) URL: #1423
rapidsai · Mar 30, 2021 · e60d9f7 · e60d9f7
1 parent ce80798
commit e60d9f7
Show file tree

Hide file tree

Showing 22 changed files with 1,415 additions and 1,557 deletions.
diff --git a/cpp/include/algorithms.hpp b/cpp/include/algorithms.hpp
@@ -14,10 +14,14 @@
  * limitations under the License.
  */
 #pragma once
+
+#include <dendrogram.hpp>
 #include <experimental/graph.hpp>
 #include <experimental/graph_view.hpp>
+
 #include <graph.hpp>
 #include <internals.hpp>
+
 #include <raft/handle.hpp>
 
 namespace cugraph {
@@ -612,7 +616,7 @@ weight_t hungarian(raft::handle_t const &handle,
  *
  * @throws     cugraph::logic_error when an error occurs.
  *
- * @tparam     graph_t               Type of graph
+ * @tparam     graph_view_t          Type of graph
  *
  * @param[in]  handle                Library handle (RAFT). If a communicator is set in the handle,
  * @param[in]  graph                 input graph object (CSR)
@@ -629,13 +633,74 @@ weight_t hungarian(raft::handle_t const &handle,
  *                                     2) modularity of the returned clustering
  *
  */
-template <typename graph_t>
-std::pair<size_t, typename graph_t::weight_type> louvain(
+template <typename graph_view_t>
+std::pair<size_t, typename graph_view_t::weight_type> louvain(
   raft::handle_t const &handle,
-  graph_t const &graph,
-  typename graph_t::vertex_type *clustering,
-  size_t max_level                         = 100,
-  typename graph_t::weight_type resolution = typename graph_t::weight_type{1});
+  graph_view_t const &graph_view,
+  typename graph_view_t::vertex_type *clustering,
+  size_t max_level                              = 100,
+  typename graph_view_t::weight_type resolution = typename graph_view_t::weight_type{1});
+
+/**
+ * @brief      Louvain implementation, returning dendrogram
+ *
+ * Compute a clustering of the graph by maximizing modularity
+ *
+ * Computed using the Louvain method described in:
+ *
+ *    VD Blondel, J-L Guillaume, R Lambiotte and E Lefebvre: Fast unfolding of
+ *    community hierarchies in large networks, J Stat Mech P10008 (2008),
+ *    http://arxiv.org/abs/0803.0476
+ *
+ * @throws     cugraph::logic_error when an error occurs.
+ *
+ * @tparam     graph_view_t          Type of graph
+ *
+ * @param[in]  handle                Library handle (RAFT)
+ * @param[in]  graph_view            Input graph view object (CSR)
+ * @param[in]  max_level             (optional) maximum number of levels to run (default 100)
+ * @param[in]  resolution            (optional) The value of the resolution parameter to use.
+ *                                   Called gamma in the modularity formula, this changes the size
+ *                                   of the communities.  Higher resolutions lead to more smaller
+ *                                   communities, lower resolutions lead to fewer larger
+ *                                   communities. (default 1)
+ *
+ * @return                           a pair containing:
+ *                                     1) unique pointer to dendrogram
+ *                                     2) modularity of the returned clustering
+ *
+ */
+template <typename graph_view_t>
+std::pair<std::unique_ptr<Dendrogram<typename graph_view_t::vertex_type>>,
+          typename graph_view_t::weight_type>
+louvain(raft::handle_t const &handle,
+        graph_view_t const &graph_view,
+        size_t max_level                              = 100,
+        typename graph_view_t::weight_type resolution = typename graph_view_t::weight_type{1});
+
+/**
+ * @brief      Flatten a Dendrogram at a particular level
+ *
+ * A Dendrogram represents a hierarchical clustering/partitioning of
+ * a graph.  This function will flatten the hierarchical clustering into
+ * a label for each vertex representing the final cluster/partition to
+ * which it is assigned
+ *
+ * @throws     cugraph::logic_error when an error occurs.
+ *
+ * @tparam     graph_view_t          Type of graph
+ *
+ * @param[in]  handle                Library handle (RAFT). If a communicator is set in the handle,
+ * @param[in]  graph                 input graph object
+ * @param[in]  dendrogram            input dendrogram object
+ * @param[out] clustering            Pointer to device array where the clustering should be stored
+ *
+ */
+template <typename graph_view_t>
+void flatten_dendrogram(raft::handle_t const &handle,
+                        graph_view_t const &graph_view,
+                        Dendrogram<typename graph_view_t::vertex_type> const &dendrogram,
+                        typename graph_view_t::vertex_type *clustering);
 
 /**
  * @brief      Leiden implementation

diff --git a/cpp/src/community/dendrogram.cuh → cpp/include/dendrogram.hpp b/cpp/src/community/dendrogram.cuh → cpp/include/dendrogram.hpp
@@ -15,7 +15,7 @@
  */
 #pragma once
 
-#include <rmm/device_buffer.hpp>
+#include <rmm/device_uvector.hpp>
 
 #include <memory>
 #include <vector>
@@ -25,30 +25,26 @@ namespace cugraph {
 template <typename vertex_t>
 class Dendrogram {
  public:
-  void add_level(vertex_t num_verts,
+  void add_level(vertex_t first_index,
+                 vertex_t num_verts,
                  cudaStream_t stream                 = 0,
                  rmm::mr::device_memory_resource *mr = rmm::mr::get_current_device_resource())
   {
-    level_ptr_.push_back(
-      std::make_unique<rmm::device_buffer>(num_verts * sizeof(vertex_t), stream, mr));
-    level_size_.push_back(num_verts);
+    level_ptr_.push_back(std::make_unique<rmm::device_uvector<vertex_t>>(num_verts, stream, mr));
+    level_first_index_.push_back(first_index);
   }
 
-  size_t current_level() const { return level_size_.size() - 1; }
+  size_t current_level() const { return level_ptr_.size() - 1; }
 
-  size_t num_levels() const { return level_size_.size(); }
+  size_t num_levels() const { return level_ptr_.size(); }
 
-  vertex_t const *get_level_ptr_nocheck(size_t level) const
-  {
-    return static_cast<vertex_t const *>(level_ptr_[level]->data());
-  }
+  vertex_t const *get_level_ptr_nocheck(size_t level) const { return level_ptr_[level]->data(); }
 
-  vertex_t *get_level_ptr_nocheck(size_t level)
-  {
-    return static_cast<vertex_t *>(level_ptr_[level]->data());
-  }
+  vertex_t *get_level_ptr_nocheck(size_t level) { return level_ptr_[level]->data(); }
 
-  vertex_t get_level_size_nocheck(size_t level) const { return level_size_[level]; }
+  size_t get_level_size_nocheck(size_t level) const { return level_ptr_[level]->size(); }
+
+  vertex_t get_level_first_index_nocheck(size_t level) const { return level_first_index_[level]; }
 
   vertex_t const *current_level_begin() const { return get_level_ptr_nocheck(current_level()); }
 
@@ -58,11 +54,16 @@ class Dendrogram {
 
   vertex_t *current_level_end() { return current_level_begin() + current_level_size(); }
 
-  vertex_t current_level_size() const { return get_level_size_nocheck(current_level()); }
+  size_t current_level_size() const { return get_level_size_nocheck(current_level()); }
+
+  vertex_t current_level_first_index() const
+  {
+    return get_level_first_index_nocheck(current_level());
+  }
 
  private:
-  std::vector<vertex_t> level_size_;
-  std::vector<std::unique_ptr<rmm::device_buffer>> level_ptr_;
+  std::vector<vertex_t> level_first_index_;
+  std::vector<std::unique_ptr<rmm::device_uvector<vertex_t>>> level_ptr_;
 };
 
 }  // namespace cugraph
diff --git a/.../experimental/include_cuco_static_map.cuh → .../experimental/include_cuco_static_map.cuh b/.../experimental/include_cuco_static_map.cuh → .../experimental/include_cuco_static_map.cuh
diff --git a/cpp/include/graph.hpp b/cpp/include/graph.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -69,6 +69,10 @@ class GraphViewBase {
   edge_t *local_edges;
   vertex_t *local_offsets;
 
+  vertex_t get_number_of_vertices() const { return number_of_vertices; }
+
+  vertex_t get_local_vertex_first() const { return vertex_t{0}; }
+
   /**
    * @brief      Fill the identifiers array with the vertex identifiers.
    *

diff --git a/cpp/include/patterns/copy_v_transform_reduce_key_aggregated_out_nbr.cuh b/cpp/include/patterns/copy_v_transform_reduce_key_aggregated_out_nbr.cuh
@@ -27,7 +27,7 @@
 
 #include <raft/handle.hpp>
 
-#include <cuco/static_map.cuh>
+#include <experimental/include_cuco_static_map.cuh>
 
 #include <type_traits>
 

diff --git a/cpp/include/patterns/transform_reduce_by_adj_matrix_row_col_key_e.cuh b/cpp/include/patterns/transform_reduce_by_adj_matrix_row_col_key_e.cuh
@@ -25,8 +25,6 @@
 
 #include <raft/handle.hpp>
 
-#include <cuco/static_map.cuh>
-
 #include <type_traits>
 
 namespace cugraph {

diff --git a/cpp/include/utilities/collect_comm.cuh b/cpp/include/utilities/collect_comm.cuh
@@ -16,6 +16,7 @@
 #pragma once
 
 #include <experimental/graph.hpp>
+#include <experimental/include_cuco_static_map.cuh>
 #include <utilities/dataframe_buffer.cuh>
 #include <utilities/shuffle_comm.cuh>
 
@@ -24,7 +25,6 @@
 #include <rmm/exec_policy.hpp>
 
 #include <thrust/distance.h>
-#include <cuco/static_map.cuh>
 
 #include <iterator>
 #include <memory>

diff --git a/cpp/include/utilities/device_comm.cuh b/cpp/include/utilities/device_comm.cuh
@@ -238,10 +238,12 @@ template <typename InputIterator, typename OutputIterator, size_t I>
 struct device_sendrecv_tuple_iterator_element_impl<InputIterator, OutputIterator, I, I> {
   void run(raft::comms::comms_t const& comm,
            InputIterator input_first,
-           size_t count,
+           size_t tx_count,
            int dst,
-           int base_tag,
-           raft::comms::request_t* requests) const
+           OutputIterator output_first,
+           size_t rx_count,
+           int src,
+           cudaStream_t stream) const
   {
   }
 };
@@ -460,7 +462,7 @@ struct device_reduce_tuple_iterator_element_impl {
                        op,
                        root,
                        stream);
-    device_reduce_tuple_iterator_element_impl<InputIterator, OutputIterator, I + 1, N>(
+    device_reduce_tuple_iterator_element_impl<InputIterator, OutputIterator, I + 1, N>().run(
       comm, input_first, output_first, count, op, root, stream);
   }
 };
@@ -889,9 +891,11 @@ device_reduce(raft::comms::comms_t const& comm,
   size_t constexpr tuple_size =
     thrust::tuple_size<typename thrust::iterator_traits<InputIterator>::value_type>::value;
 
-  detail::
-    device_reduce_tuple_iterator_element_impl<InputIterator, OutputIterator, size_t{0}, tuple_size>(
-      comm, input_first, output_first, count, op, root, stream);
+  detail::device_reduce_tuple_iterator_element_impl<InputIterator,
+                                                    OutputIterator,
+                                                    size_t{0},
+                                                    tuple_size>()
+    .run(comm, input_first, output_first, count, op, root, stream);
 }
 
 template <typename InputIterator, typename OutputIterator>

diff --git a/cpp/src/community/ecg.cu b/cpp/src/community/ecg.cu
@@ -117,7 +117,7 @@ class EcgLouvain : public cugraph::Louvain<graph_type> {
 
   void initialize_dendrogram_level(vertex_t num_vertices) override
   {
-    this->dendrogram_->add_level(num_vertices);
+    this->dendrogram_->add_level(0, num_vertices);
 
     get_permutation_vector(
       num_vertices, seed_, this->dendrogram_->current_level_begin(), this->stream_);

diff --git a/cpp/src/community/flatten_dendrogram.cuh b/cpp/src/community/flatten_dendrogram.cuh
@@ -15,7 +15,7 @@
  */
 #pragma once
 
-#include <community/dendrogram.cuh>
+#include <dendrogram.hpp>
 #include <experimental/graph_functions.hpp>
 
 #include <rmm/thrust_rmm_allocator.h>
@@ -31,23 +31,28 @@ void partition_at_level(raft::handle_t const &handle,
                         size_t level)
 {
   vertex_t local_num_verts = dendrogram.get_level_size_nocheck(0);
+  rmm::device_uvector<vertex_t> local_vertex_ids_v(local_num_verts, handle.get_stream());
 
-  thrust::copy(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()),
-               d_vertex_ids,
-               d_vertex_ids + local_num_verts,
-               d_partition);
-
-  std::for_each(thrust::make_counting_iterator<size_t>(0),
-                thrust::make_counting_iterator<size_t>(level),
-                [&handle, &dendrogram, d_vertex_ids, &d_partition, local_num_verts](size_t l) {
-                  cugraph::experimental::relabel<vertex_t, multi_gpu>(
-                    handle,
-                    std::tuple<vertex_t const *, vertex_t const *>(
-                      d_vertex_ids, dendrogram.get_level_ptr_nocheck(l)),
-                    dendrogram.get_level_size_nocheck(l),
-                    d_partition,
-                    local_num_verts);
-                });
+  raft::copy(d_partition, d_vertex_ids, local_num_verts, handle.get_stream());
+
+  std::for_each(
+    thrust::make_counting_iterator<size_t>(0),
+    thrust::make_counting_iterator<size_t>(level),
+    [&handle, &dendrogram, &local_vertex_ids_v, d_vertex_ids, &d_partition, local_num_verts](
+      size_t l) {
+      thrust::sequence(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()),
+                       local_vertex_ids_v.begin(),
+                       local_vertex_ids_v.begin() + dendrogram.get_level_size_nocheck(l),
+                       dendrogram.get_level_first_index_nocheck(l));
+
+      cugraph::experimental::relabel<vertex_t, multi_gpu>(
+        handle,
+        std::tuple<vertex_t const *, vertex_t const *>(local_vertex_ids_v.data(),
+                                                       dendrogram.get_level_ptr_nocheck(l)),
+        dendrogram.get_level_size_nocheck(l),
+        d_partition,
+        local_num_verts);
+    });
 }
 
 }  // namespace cugraph
diff --git a/cpp/src/community/leiden.cuh b/cpp/src/community/leiden.cuh
@@ -132,7 +132,7 @@ class Leiden : public Louvain<graph_type> {
       //
       //  Initialize every cluster to reference each vertex to itself
       //
-      this->dendrogram_->add_level(current_graph.number_of_vertices);
+      this->dendrogram_->add_level(0, current_graph.number_of_vertices);
 
       thrust::sequence(rmm::exec_policy(this->stream_)->on(this->stream_),
                        this->dendrogram_->current_level_begin(),