Implement eigenvector centrality

.gitignore

@@ -11,6 +11,7 @@ __pycache__
 .lock
 *.swp
 *.pytest_cache
+*~
 DartConfiguration.tcl
 .DS_Store
 

cpp/CMakeLists.txt

@@ -230,6 +230,8 @@ add_library(cugraph
     src/link_analysis/pagerank_mg.cu
     src/centrality/katz_centrality_sg.cu
     src/centrality/katz_centrality_mg.cu
+    src/centrality/eigenvector_centrality_sg.cu
+    src/centrality/eigenvector_centrality_mg.cu
     src/serialization/serializer.cu
     src/tree/mst.cu
     src/components/weakly_connected_components_sg.cu

cpp/include/cugraph/algorithms.hpp

@@ -1221,6 +1221,40 @@ void pagerank(raft::handle_t const& handle,
               bool has_initial_guess  = false,
               bool do_expensive_check = false);
 
+/**
+ * @brief Compute Eigenvector Centrality scores.
+ *
+ * This function computes eigenvector centrality scores using the power method.
+ *
+ * @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 multi_gpu Flag indicating whether template instantiation should target single-GPU (false)
+ * or multi-GPU (true).
+ * @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 initial_centralities Optional device span containing initial values for the eigenvector
+ * centralities
+ * @param epsilon Error tolerance to check convergence. Convergence is assumed if the sum of the
+ * differences in eigenvector 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 power iterations.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
+ * @return device vector containing the centralities.
+ */
+template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
+rmm::device_uvector<weight_t> eigenvector_centrality(
+  raft::handle_t const& handle,
+  graph_view_t<vertex_t, edge_t, weight_t, true, multi_gpu> const& graph_view,
+  std::optional<raft::device_span<weight_t const>> initial_centralities,
+  weight_t epsilon,
+  size_t max_iterations   = 500,
+  bool do_expensive_check = false);
+
 /**
  * @brief Compute HITS scores.
  *

cpp/include/cugraph/detail/utility_wrappers.hpp

@@ -45,19 +45,6 @@ void uniform_random_fill(rmm::cuda_stream_view const& stream_view,
                          value_t max_value,
                          uint64_t seed);
 
-/**
- * @brief    Normalize the values in an array
- *
- * @tparam      value_t      type of the value to operate on
- *
- * @param[in]   stream_view  stream view
- * @param[out]  d_value      device array to reduce
- * @param[in]   size         number of elements in array
- *
- */
-template <typename value_t>
-void normalize(rmm::cuda_stream_view const& stream_view, value_t* d_value, size_t size);
-
 /**
  * @brief    Fill a buffer with a sequence of values
  *

cpp/src/c_api/eigenvector_centrality.cpp

@@ -78,22 +78,12 @@ struct eigenvector_centrality_functor : public cugraph::c_api::abstract_functor
 
       auto number_map = reinterpret_cast<rmm::device_uvector<vertex_t>*>(graph_->number_map_);
 
-      rmm::device_uvector<weight_t> centralities(graph_view.local_vertex_partition_range_size(),
-                                                 handle_.get_stream());
-
-      // FIXME:  For now we'll call pagerank which returns a similarly formatted thing
-      cugraph::pagerank<vertex_t, edge_t, weight_t, weight_t, multi_gpu>(
+      auto centralities = cugraph::eigenvector_centrality<vertex_t, edge_t, weight_t, multi_gpu>(
         handle_,
         graph_view,
-        std::nullopt,
-        std::nullopt,
-        std::nullopt,
-        std::nullopt,
-        centralities.data(),
-        weight_t{0.95},
+        std::optional<raft::device_span<weight_t>>{},
         static_cast<weight_t>(epsilon_),
         max_iterations_,
-        false,
         do_expensive_check_);
 
       rmm::device_uvector<vertex_t> vertex_ids(graph_view.local_vertex_partition_range_size(),

cpp/src/centrality/eigenvector_centrality_impl.cuh

@@ -0,0 +1,166 @@
+/*
+ * Copyright (c) 2022, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#pragma once
+
+#include <cugraph/algorithms.hpp>
+#include <cugraph/detail/utility_wrappers.hpp>
+#include <cugraph/graph_view.hpp>
+#include <cugraph/prims/count_if_e.cuh>
+#include <cugraph/prims/count_if_v.cuh>
+#include <cugraph/prims/edge_partition_src_dst_property.cuh>
+#include <cugraph/prims/per_v_transform_reduce_incoming_outgoing_e.cuh>
+#include <cugraph/prims/reduce_v.cuh>
+#include <cugraph/prims/transform_reduce_v.cuh>
+#include <cugraph/prims/update_edge_partition_src_dst_property.cuh>
+#include <cugraph/utilities/error.hpp>
+
+#include <raft/handle.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <thrust/fill.h>
+#include <thrust/iterator/constant_iterator.h>
+#include <thrust/iterator/zip_iterator.h>
+#include <thrust/transform.h>
+#include <thrust/tuple.h>
+
+namespace cugraph {
+namespace detail {
+
+template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
+rmm::device_uvector<weight_t> eigenvector_centrality(
+  raft::handle_t const& handle,
+  graph_view_t<vertex_t, edge_t, weight_t, true, multi_gpu> const& pull_graph_view,
+  std::optional<raft::device_span<weight_t const>> initial_centralities,
+  weight_t epsilon,
+  size_t max_iterations,
+  bool do_expensive_check)
+{
+  using GraphViewType     = graph_view_t<vertex_t, edge_t, weight_t, true, multi_gpu>;
+  auto const num_vertices = pull_graph_view.number_of_vertices();
+  if (num_vertices == 0) { return rmm::device_uvector<weight_t>(0, handle.get_stream()); }
+
+  if (do_expensive_check) {
+    if (pull_graph_view.is_weighted()) {
+      auto num_nonpositive_edge_weights =
+        count_if_e(handle,
+                   pull_graph_view,
+                   dummy_property_t<vertex_t>{}.device_view(),
+                   dummy_property_t<vertex_t>{}.device_view(),
+                   [] __device__(vertex_t, vertex_t, weight_t w, auto, auto) { return w <= 0.0; });
+      CUGRAPH_EXPECTS(num_nonpositive_edge_weights == 0,
+                      "Invalid input argument: input graph should have postive edge weights.");
+    }
+  }
+
+  rmm::device_uvector<weight_t> centralities(pull_graph_view.local_vertex_partition_range_size(),
+                                             handle.get_stream());
+  if (initial_centralities) {
+    thrust::copy(handle.get_thrust_policy(),
+                 initial_centralities->begin(),
+                 initial_centralities->end(),
+                 centralities.begin());
+  } else {
+    thrust::fill(handle.get_thrust_policy(),
+                 centralities.begin(),
+                 centralities.end(),
+                 weight_t{1.0} / static_cast<weight_t>(num_vertices));
+  }
+
+  // Power iteration
+  rmm::device_uvector<weight_t> old_centralities(centralities.size(), handle.get_stream());
+
+  edge_partition_src_property_t<GraphViewType, weight_t> edge_partition_src_centralities(
+    handle, pull_graph_view);
+
+  size_t iter{0};
+  while (true) {
+    thrust::copy(handle.get_thrust_policy(),
+                 centralities.begin(),
+                 centralities.end(),
+                 old_centralities.data());
+
+    update_edge_partition_src_property(
+      handle, pull_graph_view, centralities.begin(), edge_partition_src_centralities);
+
+    per_v_transform_reduce_incoming_e(
+      handle,
+      pull_graph_view,
+      edge_partition_src_centralities.device_view(),
+      dummy_property_t<vertex_t>{}.device_view(),
+      [] __device__(vertex_t, vertex_t, weight_t w, auto src_val, auto) { return src_val * w; },
+      weight_t{0},
+      centralities.begin());
+
+    // Normalize the centralities
+    auto hypotenuse = sqrt(transform_reduce_v(
+      handle,
+      pull_graph_view,
+      centralities.begin(),
+      [] __device__(auto, auto val) { return val * val; },
+      weight_t{0.0}));
+
+    thrust::transform(handle.get_thrust_policy(),
+                      centralities.begin(),
+                      centralities.end(),
+                      centralities.begin(),
+                      [hypotenuse] __device__(auto val) { return val / hypotenuse; });
+
+    auto diff_sum = transform_reduce_v(
+      handle,
+      pull_graph_view,
+      thrust::make_zip_iterator(thrust::make_tuple(centralities.begin(), old_centralities.data())),
+      [] __device__(auto, auto val) { return std::abs(thrust::get<0>(val) - thrust::get<1>(val)); },
+      weight_t{0.0});
+
+    iter++;
+
+    if (diff_sum < (pull_graph_view.number_of_vertices() * epsilon)) {
+      break;
+    } else if (iter >= max_iterations) {
+      CUGRAPH_FAIL("Eigenvector Centrality failed to converge.");
+    }
+  }
+
+  return centralities;
+}
+
+}  // namespace detail
+
+template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
+rmm::device_uvector<weight_t> eigenvector_centrality(
+  raft::handle_t const& handle,
+  graph_view_t<vertex_t, edge_t, weight_t, true, multi_gpu> const& graph_view,
+  std::optional<raft::device_span<weight_t const>> initial_centralities,
+  weight_t epsilon,
+  size_t max_iterations,
+  bool do_expensive_check)
+{
+  static_assert(std::is_integral<vertex_t>::value,
+                "GraphViewType::vertex_type should be integral.");
+  static_assert(std::is_floating_point<weight_t>::value,
+                "weight_t should be a floating-point type.");
+
+  CUGRAPH_EXPECTS(epsilon >= 0.0, "Invalid input argument: epsilon should be non-negative.");
+  if (initial_centralities)
+    CUGRAPH_EXPECTS(initial_centralities->size() ==
+                      static_cast<size_t>(graph_view.local_vertex_partition_range_size()),
+                    "Centralities should be same size as vertex range");
+
+  return detail::eigenvector_centrality(
+    handle, graph_view, initial_centralities, epsilon, max_iterations, do_expensive_check);
+}
+
+}  // namespace cugraph

cpp/src/centrality/eigenvector_centrality_mg.cu

@@ -0,0 +1,69 @@
+/*
+ * Copyright (c) 2022, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#include <centrality/eigenvector_centrality_impl.cuh>
+
+namespace cugraph {
+
+// MG instantiation
+template rmm::device_uvector<float> eigenvector_centrality(
+  raft::handle_t const& handle,
+  graph_view_t<int32_t, int32_t, float, true, true> const& graph_view,
+  std::optional<raft::device_span<float const>> initial_centralities,
+  float epsilon,
+  size_t max_iterations,
+  bool do_expensive_check);
+
+template rmm::device_uvector<float> eigenvector_centrality(
+  raft::handle_t const& handle,
+  graph_view_t<int32_t, int64_t, float, true, true> const& graph_view,
+  std::optional<raft::device_span<float const>> initial_centralities,
+  float epsilon,
+  size_t max_iterations,
+  bool do_expensive_check);
+
+template rmm::device_uvector<float> eigenvector_centrality(
+  raft::handle_t const& handle,
+  graph_view_t<int64_t, int64_t, float, true, true> const& graph_view,
+  std::optional<raft::device_span<float const>> initial_centralities,
+  float epsilon,
+  size_t max_iterations,
+  bool do_expensive_check);
+
+template rmm::device_uvector<double> eigenvector_centrality(
+  raft::handle_t const& handle,
+  graph_view_t<int32_t, int32_t, double, true, true> const& graph_view,
+  std::optional<raft::device_span<double const>> initial_centralities,
+  double epsilon,
+  size_t max_iterations,
+  bool do_expensive_check);
+
+template rmm::device_uvector<double> eigenvector_centrality(
+  raft::handle_t const& handle,
+  graph_view_t<int32_t, int64_t, double, true, true> const& graph_view,
+  std::optional<raft::device_span<double const>> initial_centralities,
+  double epsilon,
+  size_t max_iterations,
+  bool do_expensive_check);
+
+template rmm::device_uvector<double> eigenvector_centrality(
+  raft::handle_t const& handle,
+  graph_view_t<int64_t, int64_t, double, true, true> const& graph_view,
+  std::optional<raft::device_span<double const>> initial_centralities,
+  double epsilon,
+  size_t max_iterations,
+  bool do_expensive_check);
+
+}  // namespace cugraph