Update Louvain to use new graph primitives and pattern accelerators

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 {
@@ -637,6 +641,66 @@ std::pair<size_t, typename graph_t::weight_type> louvain(
   size_t max_level                         = 100,
   typename graph_t::weight_type resolution = typename graph_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_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)
+ * @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_t>
+std::pair<std::unique_ptr<Dendrogram<typename graph_t::vertex_type>>, typename graph_t::weight_type>
+louvain(raft::handle_t const &handle,
+        graph_t const &graph,
+        size_t max_level                         = 100,
+        typename graph_t::weight_type resolution = typename graph_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_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_t>
+void flatten_dendrogram(raft::handle_t const &handle,
+                        graph_t const &graph_view,
+                        Dendrogram<typename graph_t::vertex_type> const &dendrogram,
+                        typename graph_t::vertex_type *clustering);
+
 /**
  * @brief      Leiden implementation
  *

cpp/src/community/dendrogram.cuh → cpp/include/dendrogram.hpp

@@ -16,6 +16,7 @@
 #pragma once
 
 #include <rmm/device_buffer.hpp>
+#include <rmm/device_uvector.hpp>
 
 #include <memory>
 #include <vector>
@@ -25,30 +26,33 @@ 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());
+    //return static_cast<vertex_t const *>(level_ptr_[level]->data());
+    return level_ptr_[level]->data();
   }
 
   vertex_t *get_level_ptr_nocheck(size_t level)
   {
-    return static_cast<vertex_t *>(level_ptr_[level]->data());
+    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 +62,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

cpp/include/graph.hpp

@@ -69,6 +69,14 @@ 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.
    *

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);
   }
 };
@@ -890,7 +892,7 @@ device_reduce(raft::comms::comms_t const& comm,
     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>(
+    device_reduce_tuple_iterator_element_impl<InputIterator, OutputIterator, size_t{0}, tuple_size>().run(
       comm, input_first, output_first, count, op, root, stream);
 }
 

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_);

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

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(),