[REVIEW] Pattern accelerator based implementation of PageRank, Katz Centrality, BFS, & SSSP

CHANGELOG.md

@@ -3,6 +3,7 @@
 ## New Features
 - PR #1098 Add new graph classes to support 2D partitioning
 - PR #1124 Sub-communicator initialization for 2D partitioning support
+- PR #838 Add pattern accelerator API functions and pattern accelerator API based implementations of PageRank, Katz Centrality, BFS, and SSSP
 
 ## Improvements
 - PR 1081 MNMG Renumbering - sort partitions by degree

cpp/CMakeLists.txt

@@ -337,6 +337,10 @@ add_library(cugraph SHARED
     src/centrality/betweenness_centrality.cu
     src/experimental/graph.cu
     src/experimental/graph_view.cu
+    src/experimental/bfs.cu
+    src/experimental/sssp.cu
+    src/experimental/pagerank.cu
+    src/experimental/katz_centrality.cu
 )
 
 #

cpp/include/algorithms.hpp

@@ -15,6 +15,7 @@
  */
 #pragma once
 
+#include <experimental/graph_view.hpp>
 #include <graph.hpp>
 #include <internals.hpp>
 #include <raft/handle.hpp>
@@ -936,4 +937,178 @@ void hits(GraphCSRView<VT, ET, WT> const &graph,
 
 }  // namespace gunrock
 
+namespace experimental {
+
+/**
+ * @brief Run breadth-first search to find the distances (and predecessors) from the source
+ * vertex.
+ *
+ * This function computes the distances (minimum number of hops to reach the vertex) from the source
+ * vertex. If @p predecessors is not `nullptr`, this function calculates the predecessor of each
+ * vertex (parent vertex in the breadth-first search tree) as well.
+ *
+ * @throws cugraph::logic_error on erroneous input arguments.
+ *
+ * @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 weights. Needs to be a floating point 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 graph_view Graph view object.
+ * @param distances Pointer to the output distance array.
+ * @param predecessors Pointer to the output predecessor array or `nullptr`.
+ * @param source_vertex Source vertex to start breadth-first search (root vertex of the breath-first
+ * search tree).
+ * @param direction_optimizing If set to true, this algorithm switches between the push based
+ * breadth-first search and pull based breadth-first search depending on the size of the
+ * breadth-first search frontier (currently unsupported). This option is valid only for symmetric
+ * input graphs.
+ * @param depth_limit Sets the maximum number of breadth-first search iterations. Any vertices
+ * farther than @p depth_limit hops from @p source_vertex will be marked as unreachable.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
+ */
+template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
+void bfs(raft::handle_t &handle,
+         graph_view_t<vertex_t, edge_t, weight_t, false, multi_gpu> const &graph_view,
+         vertex_t *distances,
+         vertex_t *predecessors,
+         vertex_t source_vertex,
+         bool direction_optimizing = false,
+         vertex_t depth_limit      = std::numeric_limits<vertex_t>::max(),
+         bool do_expensive_check   = false);
+
+/**
+ * @brief Run single-source shortest-path to compute the minimum distances (and predecessors) from
+ * the source vertex.
+ *
+ * This function computes the distances (minimum edge weight sums) from the source vertex. If @p
+ * predecessors is not `nullptr`, this function calculates the predecessor of each vertex in the
+ * shortest-path as well. Graph edge weights should be non-negative.
+ *
+ * @throws cugraph::logic_error on erroneous input arguments.
+ *
+ * @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 weights. Needs to be a floating point 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 graph_view Graph view object.
+ * @param distances Pointer to the output distance array.
+ * @param predecessors Pointer to the output predecessor array or `nullptr`.
+ * @param source_vertex Source vertex to start single-source shortest-path.
+ * @param cutoff Single-source shortest-path terminates if no more vertices are reachable within the
+ * distance of @p cutoff. Any vertex farther than @p cutoff will be marked as unreachable.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
+ */
+template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
+void sssp(raft::handle_t &handle,
+          graph_view_t<vertex_t, edge_t, weight_t, false, multi_gpu> const &graph_view,
+          weight_t *distances,
+          vertex_t *predecessors,
+          vertex_t source_vertex,
+          weight_t cutoff         = std::numeric_limits<weight_t>::max(),
+          bool do_expensive_check = false);
+
+/**
+ * @brief Compute PageRank scores.
+ *
+ * This function computes general (if @p personalization_vertices is `nullptr`) or personalized (if
+ * @p personalization_vertices is not `nullptr`.) PageRank scores.
+ *
+ * @throws cugraph::logic_error on erroneous input arguments or if fails to converge before @p
+ * max_iterations.
+ *
+ * @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 weights. Needs to be a floating point type.
+ * @tparam result_t Type of PageRank scores.
+ * @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 adj_matrix_row_out_weight_sums Pointer to an array storing sums of out-going edge weights
+ * for the vertices in the rows of the graph adjacency matrix (for re-use) or `nullptr`. If
+ * `nullptr`, these values are freshly computed. Computing these values outsid this function reduces
+ * the number of memoray allocations/deallocations and computing if a user repeatedly computes
+ * PageRank scores using the same graph with different personalization vectors.
+ * @param personalization_vertices Pointer to an array storing personalization vertex identifiers
+ * (compute personalized PageRank) or `nullptr` (compute general PageRank).
+ * @param personalization_values Pointer to an array storing personalization values for the vertices
+ * in the personalization set. Relevant only if @p personalization_vertices is not `nullptr`.
+ * @param personalization_vector_size Size of the personalization set. If @personalization_vertices
+ * is not `nullptr`, the sizes of the arrays pointed by @p personalization_vertices and @p
+ * personalization_values should be @p personalization_vector_size.
+ * @param pageranks Pointer to the output PageRank score array.
+ * @param alpha PageRank damping factor.
+ * @param epsilon Error tolerance to check convergence. Convergence is assumed if the sum of the
+ * differences in PageRank values between two consecutive iterations is less than the number of
+ * vertices in the graph multiplied by @p epsilon.
+ * @param max_iterations Maximum number of PageRank iterations.
+ * @param has_initial_guess If set to `true`, values in the PageRank output array (pointed by @p
+ * pageranks) is used as initial PageRank values. If false, initial PageRank values are set to 1.0
+ * divided by the number of vertices in the graph.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
+ */
+template <typename vertex_t, typename edge_t, typename weight_t, typename result_t, bool multi_gpu>
+void pagerank(raft::handle_t &handle,
+              graph_view_t<vertex_t, edge_t, weight_t, true, multi_gpu> const &graph_view,
+              weight_t *adj_matrix_row_out_weight_sums,
+              vertex_t *personalization_vertices,
+              result_t *personalization_values,
+              vertex_t personalization_vector_size,
+              result_t *pageranks,
+              result_t alpha,
+              result_t epsilon,
+              size_t max_iterations   = 500,
+              bool has_initial_guess  = false,
+              bool do_expensive_check = false);
+
+/**
+ * @brief Compute Katz Centrality scores.
+ *
+ * This function computes Katz Centrality scores.
+ *
+ * @throws cugraph::logic_error on erroneous input arguments or if fails to converge before @p
+ * max_iterations.
+ *
+ * @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 weights. Needs to be a floating point type.
+ * @tparam result_t Type of Katz Centrality scores.
+ * @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 betas Pointer to an array holding the values to be added to each vertex's new Katz
+ * Centrality score in every iteration or `nullptr`. If set to `nullptr`, constant @p beta is used
+ * instead.
+ * @param katz_centralities Pointer to the output Katz Centrality score array.
+ * @param alpha Katz Centrality attenuation factor. This should be smaller than the inverse of the
+ * maximum eigenvalue of the adjacency matrix of @p graph.
+ * @param beta Constant value to be added to each vertex's new Katz Centrality score in every
+ * iteration. Relevant only when @p betas is `nullptr`.
+ * @param epsilon Error tolerance to check convergence. Convergence is assuemd if the sum of the
+ * differences in Katz Centrality values between two consecutive iterations is less than the number
+ * of vertices in the graph multiplied by @p epsilon.
+ * @param max_iterations Maximum number of Katz Centrality iterations.
+ * @param has_initial_guess If set to `true`, values in the Katz Centrality output array (pointed by
+ * @p katz_centralities) is used as initial Katz Centrality values. If false, zeros are used as
+ * initial Katz Centrality values.
+ * @param normalize If set to `true`, final Katz Centrality scores are normalized (the L2-norm of
+ * the returned Katz Centrality score array is 1.0) before returning.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
+ */
+template <typename vertex_t, typename edge_t, typename weight_t, typename result_t, bool multi_gpu>
+void katz_centrality(raft::handle_t &handle,
+                     graph_view_t<vertex_t, edge_t, weight_t, true, multi_gpu> const &graph_view,
+                     result_t *betas,
+                     result_t *katz_centralities,
+                     result_t alpha,
+                     result_t beta,
+                     result_t epsilon,
+                     size_t max_iterations   = 500,
+                     bool has_initial_guess  = false,
+                     bool normalize          = false,
+                     bool do_expensive_check = false);
+
+}  // namespace experimental
+
 }  // namespace cugraph

cpp/include/experimental/detail/graph_utils.cuh

@@ -16,6 +16,7 @@
 #pragma once
 
 #include <experimental/graph_view.hpp>
+#include <partition_manager.hpp>
 
 #include <rmm/thrust_rmm_allocator.h>
 #include <raft/handle.hpp>
@@ -39,10 +40,10 @@ rmm::device_uvector<edge_t> compute_major_degree(
   std::vector<edge_t const *> const &adj_matrix_partition_offsets,
   partition_t<vertex_t> const &partition)
 {
-  auto &comm_p_row           = handle.get_subcomm(comm_p_row_key);
+  auto &comm_p_row           = handle.get_subcomm(cugraph::partition_2d::key_naming_t().row_name());
   auto const comm_p_row_rank = comm_p_row.get_rank();
   auto const comm_p_row_size = comm_p_row.get_size();
-  auto &comm_p_col           = handle.get_subcomm(comm_p_col_key);
+  auto &comm_p_col           = handle.get_subcomm(cugraph::partition_2d::key_naming_t().col_name());
   auto const comm_p_col_rank = comm_p_col.get_rank();
   auto const comm_p_col_size = comm_p_col.get_size();
 
@@ -74,11 +75,11 @@ rmm::device_uvector<edge_t> compute_major_degree(
     vertex_t major_first{};
     vertex_t major_last{};
     std::tie(major_first, major_last) = partition.get_vertex_partition_range(vertex_partition_idx);
-    auto p_offsets                    = partition.is_hypergraph_partitioned()
-                       ? adj_matrix_partition_offsets[i]
-                       : adj_matrix_partition_offsets[0] +
-                           (major_first - partition.get_vertex_partition_range_first(
-                                            comm_p_col_size * comm_p_row_rank));
+    auto p_offsets =
+      partition.is_hypergraph_partitioned()
+        ? adj_matrix_partition_offsets[i]
+        : adj_matrix_partition_offsets[0] +
+            (major_first - partition.get_vertex_partition_first(comm_p_col_size * comm_p_row_rank));
     thrust::transform(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()),
                       thrust::make_counting_iterator(vertex_t{0}),
                       thrust::make_counting_iterator(major_last - major_first),

cpp/include/experimental/graph.hpp

@@ -72,8 +72,7 @@ class graph_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu, std::enab
 
   vertex_t get_number_of_local_vertices() const
   {
-    return partition_.get_vertex_partition_range_last() -
-           partition_.get_vertex_partition_range_first();
+    return partition_.get_vertex_partition_last() - partition_.get_vertex_partition_first();
   }
 
   graph_view_t<vertex_t, edge_t, weight_t, store_transposed, multi_gpu> view()