HITS primitive based implementation

cpp/CMakeLists.txt

@@ -163,7 +163,6 @@ add_library(cugraph SHARED
     src/utilities/graph_bcast.cu
     src/structure/legacy/graph.cu
     src/linear_assignment/hungarian.cu
-    #src/link_analysis/gunrock_hits.cpp
     src/traversal/legacy/bfs.cu
     src/traversal/legacy/sssp.cu
     src/link_prediction/jaccard.cu
@@ -211,6 +210,8 @@ add_library(cugraph SHARED
     src/traversal/bfs_mg.cu
     src/traversal/sssp_sg.cu
     src/traversal/sssp_mg.cu
+    src/link_analysis/hits_sg.cu
+    src/link_analysis/hits_mg.cu
     src/link_analysis/pagerank_sg.cu
     src/link_analysis/pagerank_mg.cu
     src/centrality/katz_centrality_sg.cu

cpp/include/cugraph/algorithms.hpp

@@ -1256,6 +1256,47 @@ void pagerank(raft::handle_t const& handle,
               bool has_initial_guess  = false,
               bool do_expensive_check = false);
 
+/**
+ * @brief Compute HITS scores.
+ *
+ * This function computes HITS scores for the vertices of a graph
+ *
+ * @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.
+ * @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 hubs Pointer to the input/output hub score array.
+ * @param authorities Pointer to the output authorities score array.
+ * @param epsilon Error tolerance to check convergence. Convergence is assumed if the sum of the
+ * differences in hub values between two consecutive iterations is less than @p epsilon
+ * @param max_iterations Maximum number of HITS iterations.
+ * @param has_initial_guess If set to `true`, values in the hubs output array (pointed by @p
+ * hubs) is used as initial hub values. If false, initial hub values are set to 1.0
+ * divided by the number of vertices in the graph.
+ * @param normalize If set to `true`, final hub and authority scores are normalized (the L1-norm of
+ * the returned hub and authority score arrays is 1.0) before returning.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
+ * @return std::tuple<weight_t, size_t> A tuple of sum of the differences of hub scores of the last
+ * two iterations and the total number of iterations taken to reach the final result
+ */
+template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
+std::tuple<weight_t, size_t> hits(
+  raft::handle_t const& handle,
+  graph_view_t<vertex_t, edge_t, weight_t, true, multi_gpu> const& graph_view,
+  weight_t* hubs,
+  weight_t* authorities,
+  weight_t epsilon,
+  size_t max_iterations,
+  bool has_initial_hubs_guess,
+  bool normalize,
+  bool do_expensive_check);
+
 /**
  * @brief Compute Katz Centrality scores.
  *

cpp/include/cugraph/prims/property_op_utils.cuh

@@ -154,15 +154,11 @@ constexpr auto op_dispatch(raft::comms::op_t op, F&& f)
     case raft::comms::op_t::SUM: {
       return std::invoke(f, property_op<T, thrust::plus>());
     } break;
-    case raft::comms::op_t::PROD: {
-      CUGRAPH_FAIL("raft::comms::op_t::PROD is not supported for op_dispatch");
-      return std::invoke_result_t<F, property_op<T, thrust::multiplies>>{};
-    } break;
     case raft::comms::op_t::MIN: {
-      return std::invoke(f, property_op<T, thrust::less>());
+      return std::invoke(f, property_op<T, thrust::minimum>());
     } break;
     case raft::comms::op_t::MAX: {
-      return std::invoke(f, property_op<T, thrust::greater>());
+      return std::invoke(f, property_op<T, thrust::maximum>());
     } break;
     default: {
       CUGRAPH_FAIL("Unhandled raft::comms::op_t");
@@ -171,6 +167,26 @@ constexpr auto op_dispatch(raft::comms::op_t op, F&& f)
   };
 }
 
+template <typename T>
+T identity_element(raft::comms::op_t op)
+{
+  switch (op) {
+    case raft::comms::op_t::SUM: {
+      return T{0};
+    } break;
+    case raft::comms::op_t::MIN: {
+      return std::numeric_limits<T>::max();
+    } break;
+    case raft::comms::op_t::MAX: {
+      return std::numeric_limits<T>::lowest();
+    } break;
+    default: {
+      CUGRAPH_FAIL("Unhandled raft::comms::op_t");
+      return T{0};
+    }
+  };
+}
+
 template <typename Iterator, typename T>
 __device__ std::enable_if_t<thrust::detail::is_discard_iterator<Iterator>::value, void>
 atomic_accumulate_edge_op_result(Iterator iter, T const& value)

cpp/include/cugraph/prims/reduce_v.cuh

@@ -49,17 +49,19 @@ template <typename GraphViewType, typename VertexValueInputIterator, typename T>
 T reduce_v(raft::handle_t const& handle,
            GraphViewType const& graph_view,
            VertexValueInputIterator vertex_value_input_first,
-           T init               = T{},
+           T init,
            raft::comms::op_t op = raft::comms::op_t::SUM)
 {
-  auto ret = op_dispatch<T>(op, [&handle, &graph_view, vertex_value_input_first, init](auto op) {
-    return thrust::reduce(
-      handle.get_thrust_policy(),
-      vertex_value_input_first,
-      vertex_value_input_first + graph_view.get_number_of_local_vertices(),
-      ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() != 0)) ? T{} : init,
-      op);
-  });
+  auto id = identity_element<T>(op);
+  auto ret =
+    op_dispatch<T>(op, [&handle, &graph_view, vertex_value_input_first, id, init](auto op) {
+      return thrust::reduce(
+        handle.get_thrust_policy(),
+        vertex_value_input_first,
+        vertex_value_input_first + graph_view.get_number_of_local_vertices(),
+        ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() != 0)) ? id : init,
+        op);
+    });
   if constexpr (GraphViewType::is_multi_gpu) {
     ret = host_scalar_allreduce(handle.get_comms(), ret, op, handle.get_stream());
   }

cpp/include/cugraph/prims/transform_reduce_v.cuh

@@ -53,17 +53,18 @@ T transform_reduce_v(raft::handle_t const& handle,
                      GraphViewType const& graph_view,
                      VertexValueInputIterator vertex_value_input_first,
                      VertexOp v_op,
-                     T init               = T{},
+                     T init,
                      raft::comms::op_t op = raft::comms::op_t::SUM)
 {
+  auto id = identity_element<T>(op);
   auto ret =
-    op_dispatch<T>(op, [&handle, &graph_view, vertex_value_input_first, v_op, init](auto op) {
+    op_dispatch<T>(op, [&handle, &graph_view, vertex_value_input_first, v_op, id, init](auto op) {
       return thrust::transform_reduce(
         handle.get_thrust_policy(),
         vertex_value_input_first,
         vertex_value_input_first + graph_view.get_number_of_local_vertices(),
         v_op,
-        ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() != 0)) ? T{} : init,
+        ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() != 0)) ? id : init,
         op);
     });
   if (GraphViewType::is_multi_gpu) {

cpp/src/link_analysis/hits_impl.cuh

@@ -0,0 +1,203 @@
+/*
+ * Copyright (c) 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.
+ * 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/graph_view.hpp>
+#include <cugraph/prims/copy_to_adj_matrix_row_col.cuh>
+#include <cugraph/prims/copy_v_transform_reduce_in_out_nbr.cuh>
+#include <cugraph/prims/count_if_v.cuh>
+#include <cugraph/prims/reduce_v.cuh>
+#include <cugraph/prims/row_col_properties.cuh>
+#include <cugraph/prims/transform_reduce_v.cuh>
+
+#include <thrust/fill.h>
+#include <thrust/transform.h>
+
+namespace cugraph {
+namespace detail {
+template <typename GraphViewType, typename result_t>
+void normalize(raft::handle_t const& handle,
+               GraphViewType const& graph_view,
+               result_t* hubs,
+               raft::comms::op_t op)
+{
+  auto hubs_norm = reduce_v(handle,
+                            graph_view,
+                            hubs,
+                            hubs + graph_view.get_number_of_local_vertices(),
+                            identity_element<result_t>(op),
+                            op);
+  CUGRAPH_EXPECTS(hubs_norm > 0, "Norm is required to be a positive value.");
+  thrust::transform(handle.get_thrust_policy(),
+                    hubs,
+                    hubs + graph_view.get_number_of_local_vertices(),
+                    thrust::make_constant_iterator(hubs_norm),
+                    hubs,
+                    thrust::divides<result_t>());
+}
+
+template <typename GraphViewType, typename result_t>
+std::tuple<result_t, size_t> hits(raft::handle_t const& handle,
+                                  GraphViewType const& graph_view,
+                                  result_t* const hubs,
+                                  result_t* const authorities,
+                                  result_t epsilon,
+                                  size_t max_iterations,
+                                  bool has_initial_hubs_guess,
+                                  bool normalize,
+                                  bool do_expensive_check)
+{
+  using vertex_t = typename GraphViewType::vertex_type;
+  static_assert(std::is_integral<vertex_t>::value,
+                "GraphViewType::vertex_type should be integral.");
+  static_assert(std::is_floating_point<result_t>::value,
+                "result_t should be a floating-point type.");
+  static_assert(GraphViewType::is_adj_matrix_transposed,
+                "GraphViewType should support the pull model.");
+
+  auto const num_vertices = graph_view.get_number_of_vertices();
+  result_t diff_sum{std::numeric_limits<result_t>::max()};
+  size_t final_iteration_count{max_iterations};
+
+  if (num_vertices == 0) { return std::make_tuple(diff_sum, final_iteration_count); }
+
+  CUGRAPH_EXPECTS(epsilon >= 0.0, "Invalid input argument: epsilon should be non-negative.");
+
+  // Check validity of initial guess if supplied
+  if (has_initial_hubs_guess && do_expensive_check) {
+    auto num_negative_values =
+      count_if_v(handle, graph_view, hubs, [] __device__(auto val) { return val < 0.0; });
+    CUGRAPH_EXPECTS(num_negative_values == 0,
+                    "Invalid input argument: initial guess values should be non-negative.");
+  }
+
+  if (has_initial_hubs_guess) {
+    detail::normalize(handle, graph_view, hubs, raft::comms::op_t::SUM);
+  }
+
+  // Property wrappers
+  row_properties_t<GraphViewType, result_t> prev_src_hubs(handle, graph_view);
+  col_properties_t<GraphViewType, result_t> curr_dst_auth(handle, graph_view);
+  rmm::device_uvector<result_t> temp_hubs(graph_view.get_number_of_local_vertices(),
+                                          handle.get_stream());
+
+  result_t* prev_hubs = hubs;
+  result_t* curr_hubs = temp_hubs.data();
+
+  // Initialize hubs from user input if provided
+  if (has_initial_hubs_guess) {
+    copy_to_adj_matrix_row(handle, graph_view, prev_hubs, prev_src_hubs);
+  } else {
+    prev_src_hubs.fill(result_t{1.0} / num_vertices, handle.get_stream());
+    thrust::fill(handle.get_thrust_policy(),
+                 prev_hubs,
+                 prev_hubs + graph_view.get_number_of_local_vertices(),
+                 result_t{1.0} / num_vertices);
+  }
+  for (size_t iter = 0; iter < max_iterations; ++iter) {
+    // Update current destination authorities property
+    copy_v_transform_reduce_in_nbr(
+      handle,
+      graph_view,
+      prev_src_hubs.device_view(),
+      dummy_properties_t<result_t>{}.device_view(),
+      [] __device__(auto, auto, auto, auto prev_src_hub_value, auto) { return prev_src_hub_value; },
+      result_t{0},
+      authorities);
+
+    copy_to_adj_matrix_col(handle, graph_view, authorities, curr_dst_auth);
+
+    // Update current source hubs property
+    copy_v_transform_reduce_out_nbr(
+      handle,
+      graph_view,
+      dummy_properties_t<result_t>{}.device_view(),
+      curr_dst_auth.device_view(),
+      [] __device__(auto src, auto dst, auto, auto, auto curr_dst_auth_value) {
+        return curr_dst_auth_value;
+      },
+      result_t{0},
+      curr_hubs);
+
+    // Normalize current hub values
+    detail::normalize(handle, graph_view, curr_hubs, raft::comms::op_t::MAX);
+
+    // Normalize current authority values
+    detail::normalize(handle, graph_view, authorities, raft::comms::op_t::MAX);
+
+    // Test for exit condition
+    diff_sum = transform_reduce_v(
+      handle,
+      graph_view,
+      thrust::make_zip_iterator(thrust::make_tuple(curr_hubs, prev_hubs)),
+      [] __device__(auto val) { return std::abs(thrust::get<0>(val) - thrust::get<1>(val)); },
+      result_t{0});
+    if (diff_sum < epsilon) {
+      final_iteration_count = iter;
+      std::swap(prev_hubs, curr_hubs);
+      break;
+    }
+
+    copy_to_adj_matrix_row(handle, graph_view, curr_hubs, prev_src_hubs);
+
+    // Swap pointers for the next iteration
+    // After this swap call, prev_hubs has the latest value of hubs
+    std::swap(prev_hubs, curr_hubs);
+  }
+
+  if (normalize) {
+    detail::normalize(handle, graph_view, prev_hubs, raft::comms::op_t::SUM);
+    detail::normalize(handle, graph_view, authorities, raft::comms::op_t::SUM);
+  }
+
+  // Copy calculated hubs to in/out parameter if necessary
+  if (hubs != prev_hubs) {
+    thrust::copy(handle.get_thrust_policy(),
+                 prev_hubs,
+                 prev_hubs + graph_view.get_number_of_local_vertices(),
+                 hubs);
+  }
+
+  return std::make_tuple(diff_sum, final_iteration_count);
+}
+
+}  // namespace detail
+
+template <typename vertex_t, typename edge_t, typename weight_t, bool multi_gpu>
+std::tuple<weight_t, size_t> hits(
+  raft::handle_t const& handle,
+  graph_view_t<vertex_t, edge_t, weight_t, true, multi_gpu> const& graph_view,
+  weight_t* const hubs,
+  weight_t* const authorities,
+  weight_t epsilon,
+  size_t max_iterations,
+  bool has_initial_hubs_guess,
+  bool normalize,
+  bool do_expensive_check)
+{
+  return detail::hits(handle,
+                      graph_view,
+                      hubs,
+                      authorities,
+                      epsilon,
+                      max_iterations,
+                      has_initial_hubs_guess,
+                      normalize,
+                      do_expensive_check);
+}
+
+}  // namespace cugraph