Merge pull request #870 from ChuckHastings/fix_louvain

[REVIEW] Fix louvain

Author: afender

CHANGELOG.md

@@ -61,6 +61,7 @@
 - PR #852 Fix BFS Notebook
 - PR #855 Missed a file in the original SNMG PR
 - PR #860 Fix all Notebooks
+- PR #870 Fix Louvain
 - PR #889 Added missing conftest.py file to benchmarks dir
 - PR #896 opg dask infrastructure fixes
 

cpp/CMakeLists.txt

@@ -316,8 +316,9 @@ add_library(cugraph SHARED
     src/converters/renumber.cu
     src/converters/COOtoCSR.cu
     src/community/spectral_clustering.cu
+    src/community/louvain.cpp
+    src/community/louvain_kernels.cu
     src/community/ktruss.cu
-    src/community/louvain.cu
     src/community/ECG.cu
     src/community/triangles_counting.cu
     src/community/extract_subgraph_by_vertex.cu

cpp/include/algorithms.hpp

@@ -553,6 +553,62 @@ void bfs(experimental::GraphCSRView<VT, ET, WT> const &graph,
          const VT start_vertex,
          bool directed = true);
 
+/**
+ * @brief      Louvain implementation
+ *
+ * Compute a clustering of the graph by minimizing modularity
+ *
+ * @throws     cugraph::logic_error when an error occurs.
+ *
+ * @tparam VT                        Type of vertex identifiers.
+ *                                   Supported value : int (signed, 32-bit)
+ * @tparam ET                        Type of edge identifiers.
+ *                                   Supported value : int (signed, 32-bit)
+ * @tparam WT                        Type of edge weights. Supported values : float or double.
+ *
+ * @param[in]  graph                 input graph object (CSR)
+ * @param[out] final_modularity      modularity of the returned clustering
+ * @param[out] num_level             number of levels of the returned clustering
+ * @param[out] clustering            Pointer to device array where the clustering should be stored
+ * @param[in]  max_iter              (optional) maximum number of iterations to run (default 100)
+ */
+template <typename VT, typename ET, typename WT>
+void louvain(experimental::GraphCSRView<VT, ET, WT> const &graph,
+             WT *final_modularity,
+             int *num_level,
+             VT *louvain_parts,
+             int max_iter = 100);
+
+/**
+ * @brief Computes the ecg clustering of the given graph.
+ *
+ * ECG runs truncated Louvain on an ensemble of permutations of the input graph,
+ * then uses the ensemble partitions to determine weights for the input graph.
+ * The final result is found by running full Louvain on the input graph using
+ * the determined weights. See https://arxiv.org/abs/1809.05578 for further
+ * information.
+ *
+ * @throws     cugraph::logic_error when an error occurs.
+ *
+ * @tparam VT                        Type of vertex identifiers. Supported value : int (signed,
+ * 32-bit)
+ * @tparam ET                        Type of edge identifiers.  Supported value : int (signed,
+ * 32-bit)
+ * @tparam WT                        Type of edge weights. Supported values : float or double.
+ *
+ * @param[in]  graph_coo             input graph object (COO)
+ * @param[in]  graph_csr             input graph object (CSR)
+ * @param[in]  min_weight            The minimum weight parameter
+ * @param[in]  ensemble_size         The ensemble size parameter
+ * @param[out] ecg_parts             A device pointer to array where the partitioning should be
+ * written
+ */
+template <typename VT, typename ET, typename WT>
+void ecg(experimental::GraphCSRView<VT, ET, WT> const &graph_csr,
+         WT min_weight,
+         VT ensemble_size,
+         VT *ecg_parts);
+
 namespace nvgraph {
 
 /**
@@ -725,59 +781,5 @@ void analyzeClustering_ratio_cut(experimental::GraphCSRView<VT, ET, WT> const &g
                                  VT const *clustering,
                                  WT *score);
 
-/**
- * @brief      Wrapper function for Nvgraph louvain implementation
- *
- * @throws     cugraph::logic_error when an error occurs.
- *
- * @tparam VT                        Type of vertex identifiers. Supported value : int (signed,
- * 32-bit)
- * @tparam ET                        Type of edge identifiers.  Supported value : int (signed,
- * 32-bit)
- * @tparam WT                        Type of edge weights. Supported values : float or double.
- *
- * @param[in]  graph                 input graph object (CSR)
- * @param[out] final_modularity      modularity of the returned clustering
- * @param[out] num_level             number of levels of the returned clustering
- * @param[out] clustering            Pointer to device array where the clustering should be stored
- * @param[in]  max_iter              (optional) maximum number of iterations to run (default 100)
- */
-template <typename VT, typename ET, typename WT>
-void louvain(experimental::GraphCSRView<VT, ET, WT> const &graph,
-             WT *final_modularity,
-             VT *num_level,
-             VT *louvain_parts,
-             int max_iter = 100);
-
-/**
- * @brief Computes the ecg clustering of the given graph.
- *
- * ECG runs truncated Louvain on an ensemble of permutations of the input graph,
- * then uses the ensemble partitions to determine weights for the input graph.
- * The final result is found by running full Louvain on the input graph using
- * the determined weights. See https://arxiv.org/abs/1809.05578 for further
- * information.
- *
- * @throws     cugraph::logic_error when an error occurs.
- *
- * @tparam VT                        Type of vertex identifiers. Supported value : int (signed,
- * 32-bit)
- * @tparam ET                        Type of edge identifiers.  Supported value : int (signed,
- * 32-bit)
- * @tparam WT                        Type of edge weights. Supported values : float or double.
- *
- * @param[in]  graph_coo             input graph object (COO)
- * @param[in]  graph_csr             input graph object (CSR)
- * @param[in]  min_weight            The minimum weight parameter
- * @param[in]  ensemble_size         The ensemble size parameter
- * @param[out] ecg_parts             A device pointer to array where the partitioning should be
- * written
- */
-template <typename VT, typename ET, typename WT>
-void ecg(experimental::GraphCSRView<VT, ET, WT> const &graph_csr,
-         WT min_weight,
-         VT ensemble_size,
-         VT *ecg_parts);
-
 }  // namespace nvgraph
 }  // namespace cugraph

cpp/src/community/ECG.cu

@@ -92,22 +92,21 @@ struct update_functor {
  * responsible for freeing the allocated memory using ALLOC_FREE_TRY().
  */
 template <typename T>
-void get_permutation_vector(T size, T seed, T *permutation)
+void get_permutation_vector(T size, T seed, T *permutation, cudaStream_t stream)
 {
   rmm::device_vector<float> randoms_v(size);
 
   thrust::counting_iterator<uint32_t> index(seed);
   thrust::transform(
-    rmm::exec_policy(nullptr)->on(nullptr), index, index + size, randoms_v.begin(), prg());
-  thrust::sequence(rmm::exec_policy(nullptr)->on(nullptr), permutation, permutation + size, 0);
+    rmm::exec_policy(stream)->on(stream), index, index + size, randoms_v.begin(), prg());
+  thrust::sequence(rmm::exec_policy(stream)->on(stream), permutation, permutation + size, 0);
   thrust::sort_by_key(
-    rmm::exec_policy(nullptr)->on(nullptr), randoms_v.begin(), randoms_v.end(), permutation);
+    rmm::exec_policy(stream)->on(stream), randoms_v.begin(), randoms_v.end(), permutation);
 }
 
 }  // anonymous namespace
 
 namespace cugraph {
-namespace nvgraph {
 
 template <typename VT, typename ET, typename WT>
 void ecg(experimental::GraphCSRView<VT, ET, WT> const &graph,
@@ -118,7 +117,9 @@ void ecg(experimental::GraphCSRView<VT, ET, WT> const &graph,
   CUGRAPH_EXPECTS(graph.edge_data != nullptr, "API error, louvain expects a weighted graph");
   CUGRAPH_EXPECTS(ecg_parts != nullptr, "Invalid API parameter: ecg_parts is NULL");
 
-  rmm::device_vector<WT> ecg_weights_v(graph.number_of_edges, WT{0.0});
+  cudaStream_t stream{0};
+
+  rmm::device_vector<WT> ecg_weights_v(graph.edge_data, graph.edge_data + graph.number_of_edges);
 
   VT size{graph.number_of_vertices};
   VT seed{0};
@@ -133,7 +134,7 @@ void ecg(experimental::GraphCSRView<VT, ET, WT> const &graph,
     rmm::device_vector<VT> permutation_v(size);
     VT *d_permutation = permutation_v.data().get();
 
-    get_permutation_vector(size, seed, d_permutation);
+    get_permutation_vector(size, seed, d_permutation, stream);
     seed += size;
 
     detail::permute_graph<VT, ET, WT>(graph, d_permutation, permuted_graph->view());
@@ -145,25 +146,25 @@ void ecg(experimental::GraphCSRView<VT, ET, WT> const &graph,
     WT final_modularity;
     VT num_level;
 
-    cugraph::nvgraph::louvain(permuted_graph->view(), &final_modularity, &num_level, d_parts, 1);
+    cugraph::louvain(permuted_graph->view(), &final_modularity, &num_level, d_parts, 1);
 
     // For each edge in the graph determine whether the endpoints are in the same partition
     // Keep a sum for each edge of the total number of times its endpoints are in the same partition
     dim3 grid, block;
     block.x = 512;
     grid.x  = min(VT{CUDA_MAX_BLOCKS}, (graph.number_of_edges / 512 + 1));
-    match_check_kernel<<<grid, block, 0, nullptr>>>(graph.number_of_edges,
-                                                    graph.number_of_vertices,
-                                                    graph.offsets,
-                                                    graph.indices,
-                                                    permutation_v.data().get(),
-                                                    d_parts,
-                                                    ecg_weights_v.data().get());
+    match_check_kernel<<<grid, block, 0, stream>>>(graph.number_of_edges,
+                                                   graph.number_of_vertices,
+                                                   graph.offsets,
+                                                   graph.indices,
+                                                   permutation_v.data().get(),
+                                                   d_parts,
+                                                   ecg_weights_v.data().get());
   }
 
   // Set weights = min_weight + (1 - min-weight)*sum/ensemble_size
   update_functor<WT> uf(min_weight, ensemble_size);
-  thrust::transform(rmm::exec_policy(nullptr)->on(nullptr),
+  thrust::transform(rmm::exec_policy(stream)->on(stream),
                     ecg_weights_v.data().get(),
                     ecg_weights_v.data().get() + graph.number_of_edges,
                     ecg_weights_v.data().get(),
@@ -179,7 +180,7 @@ void ecg(experimental::GraphCSRView<VT, ET, WT> const &graph,
 
   WT final_modularity;
   VT num_level;
-  cugraph::nvgraph::louvain(louvain_graph, &final_modularity, &num_level, ecg_parts, 100);
+  cugraph::louvain(louvain_graph, &final_modularity, &num_level, ecg_parts, 100);
 }
 
 // Explicit template instantiations.
@@ -193,5 +194,4 @@ template void ecg<int32_t, int32_t, double>(
   double min_weight,
   int32_t ensemble_size,
   int32_t *ecg_parts);
-}  // namespace nvgraph
 }  // namespace cugraph

cpp/src/community/louvain.cu renamed to cpp/src/community/louvain.cpp

@@ -17,16 +17,20 @@
 #include <algorithms.hpp>
 #include <graph.hpp>
 
-#include <nvgraph/include/nvlouvain.cuh>
+#include <rmm/thrust_rmm_allocator.h>
+
+#include <thrust/sequence.h>
+
+#include <community/louvain_kernels.hpp>
+
 #include "utilities/error_utils.h"
 
 namespace cugraph {
-namespace nvgraph {
 
 template <typename VT, typename ET, typename WT>
 void louvain(experimental::GraphCSRView<VT, ET, WT> const &graph,
              WT *final_modularity,
-             VT *num_level,
+             int *num_level,
              VT *louvain_parts,
              int max_iter)
 {
@@ -35,38 +39,12 @@ void louvain(experimental::GraphCSRView<VT, ET, WT> const &graph,
   CUGRAPH_EXPECTS(num_level != nullptr, "API error, num_level is null");
   CUGRAPH_EXPECTS(louvain_parts != nullptr, "API error, louvain_parts is null");
 
-  std::ostream log(0);
-
-  bool weighted{true};
-
-  WT mod{0.0};
-  VT n_level{0};
-
-  nvlouvain::louvain<VT, WT>(graph.offsets,
-                             graph.indices,
-                             graph.edge_data,
-                             graph.number_of_vertices,
-                             graph.number_of_edges,
-                             weighted,
-                             false,
-                             nullptr,
-                             mod,
-                             louvain_parts,
-                             n_level,
-                             max_iter,
-                             log);
-
-  *final_modularity = mod;
-  *num_level        = n_level;
+  detail::louvain<VT, ET, WT>(graph, final_modularity, num_level, louvain_parts, max_iter);
 }
 
 template void louvain(
-  experimental::GraphCSRView<int32_t, int32_t, float> const &, float *, int32_t *, int32_t *, int);
-template void louvain(experimental::GraphCSRView<int32_t, int32_t, double> const &,
-                      double *,
-                      int32_t *,
-                      int32_t *,
-                      int);
+  experimental::GraphCSRView<int32_t, int32_t, float> const &, float *, int *, int32_t *, int);
+template void louvain(
+  experimental::GraphCSRView<int32_t, int32_t, double> const &, double *, int *, int32_t *, int);
 
-}  // namespace nvgraph
 }  // namespace cugraph