Ktruss implementation

cpp/CMakeLists.txt

@@ -199,6 +199,7 @@ set(CUGRAPH_SOURCES
     src/community/detail/refine_mg.cu
     src/community/detail/mis_sg.cu
     src/community/detail/mis_mg.cu
+    src/community/edge_triangle_count_sg.cu
     src/detail/utility_wrappers.cu
     src/structure/graph_view_mg.cu
     src/structure/remove_self_loops.cu
@@ -226,6 +227,7 @@ set(CUGRAPH_SOURCES
     src/community/legacy/ecg.cu
     src/community/egonet_sg.cu
     src/community/egonet_mg.cu
+    src/community/k_truss_sg.cu
     src/sampling/random_walks.cu
     src/sampling/random_walks_sg.cu
     src/sampling/detail/prepare_next_frontier_sg.cu

cpp/include/cugraph/algorithms.hpp

@@ -1842,7 +1842,7 @@ void weakly_connected_components(raft::handle_t const& handle,
 enum class k_core_degree_type_t { IN = 0, OUT = 1, INOUT = 2 };
 
 /**
- * @brief   Compute core numbers of individual vertices from K-core decomposition.
+ * @brief   Compute core numbers of individual vertices from K-Core decomposition.
  *
  * The input graph should not have self-loops nor multi-edges. Currently, only undirected graphs are
  * supported.
@@ -1855,11 +1855,11 @@ enum class k_core_degree_type_t { IN = 0, OUT = 1, INOUT = 2 };
  * handles to various CUDA libraries) to run graph algorithms.
  * @param graph_view Graph view object.
  * @param core_numbers Pointer to the output core number array.
- * @param degree_type Dictate whether to compute the K-core decomposition based on in-degrees,
+ * @param degree_type Dictate whether to compute the K-Core decomposition based on in-degrees,
  * out-degrees, or in-degrees + out_degrees.
- * @param k_first Find K-cores from K = k_first. Any vertices that do not belong to k_first-core
+ * @param k_first Find K-Cores from K = k_first. Any vertices that do not belong to k_first-core
  * will have core numbers of 0.
- * @param k_last Find K-cores to K = k_last. Any vertices that belong to (k_last)-core will have
+ * @param k_last Find K-Cores to K = k_last. Any vertices that belong to (k_last)-core will have
  * their core numbers set to their degrees on k_last-core.
  * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  */
@@ -1873,7 +1873,7 @@ void core_number(raft::handle_t const& handle,
                  bool do_expensive_check = false);
 
 /**
- * @brief   Extract K Core of a graph
+ * @brief   Extract K-Core of a graph
  *
  * @throws     cugraph::logic_error when an error occurs.
  *
@@ -1884,7 +1884,7 @@ void core_number(raft::handle_t const& handle,
  * @param  graph_view      Graph view object.
  * @param edge_weight_view Optional view object holding edge weights for @p graph_view.
  * @param  k               Order of the core. This value must not be negative.
- * @param degree_type Optional parameter to dictate whether to compute the K-core decomposition
+ * @param degree_type Optional parameter to dictate whether to compute the K-Core decomposition
  *                    based on in-degrees, out-degrees, or in-degrees + out_degrees.  One of @p
  *                    degree_type and @p core_numbers must be specified.
  * @param  core_numbers    Optional output from core_number algorithm.  If not specified then
@@ -2040,6 +2040,46 @@ void triangle_count(raft::handle_t const& handle,
                     raft::device_span<edge_t> counts,
                     bool do_expensive_check = false);
 
+/**
+ * @brief Count the number of triangles for each edge.
+ * @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.
+ *
+ * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
+ * handles to various CUDA libraries) to run graph algorithms.
+ * @param edgelist_srcs  List of source vertex ids
+ * @param edgelist_dsts  List of destination vertex ids
+ * @return triangle count of the edge list
+ */
+template <typename vertex_t, typename edge_t, bool store_transposed, bool multi_gpu>
+rmm::device_uvector<edge_t> edge_triangle_count(
+  raft::handle_t const& handle,
+  graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu> const& graph_view,
+  raft::device_span<vertex_t> edgelist_srcs,
+  raft::device_span<vertex_t> edgelist_dsts);
+
+/*
+ * @brief Compute K-Truss.
+ *
+ * Extract the K-Truss subgraph of a graph
+ *
+ * @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 multi_gpu Flag indicating whether template instantiation should target single-GPU (false)
+ * @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 k The desired k to be used for extracting the K-Truss subgraph
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
+ * @return edge list of the K-Truss subgraph
+ */
+template <typename vertex_t, typename edge_t, bool multi_gpu>
+std::tuple<rmm::device_uvector<vertex_t>, rmm::device_uvector<vertex_t>> k_truss(
+  raft::handle_t const& handle,
+  graph_view_t<vertex_t, edge_t, false, multi_gpu> const& graph_view,
+  edge_t k,
+  bool do_expensive_check = false);
+
 /**
  * @brief     Compute Jaccard similarity coefficient
  *

cpp/src/c_api/legacy_k_truss.cpp

@@ -81,27 +81,13 @@ struct k_truss_functor : public cugraph::c_api::abstract_functor {
       auto number_map = reinterpret_cast<rmm::device_uvector<vertex_t>*>(graph_->number_map_);
 
       auto graph_view = graph->view();
-      rmm::device_uvector<vertex_t> src(0, handle_.get_stream());
-      rmm::device_uvector<vertex_t> dst(0, handle_.get_stream());
-      std::optional<rmm::device_uvector<weight_t>> wgt{std::nullopt};
 
-      std::tie(src, dst, wgt, std::ignore) = cugraph::decompress_to_edgelist(
+      auto [result_src, result_dst] = cugraph::k_truss<vertex_t, edge_t, multi_gpu>(
         handle_,
         graph_view,
-        edge_weights ? std::make_optional(edge_weights->view()) : std::nullopt,
-        std::optional<cugraph::edge_property_view_t<edge_t, edge_t const*>>{std::nullopt},
-        std::optional<raft::device_span<vertex_t const>>(std::nullopt),
+        k_,
         do_expensive_check_);
 
-      auto [result_src, result_dst, result_wgt] = cugraph::k_truss_subgraph(
-        handle_,
-        raft::device_span<vertex_t>(src.data(), src.size()),
-        raft::device_span<vertex_t>(dst.data(), dst.size()),
-        wgt ? std::make_optional(raft::device_span<weight_t>(wgt->data(), wgt->size()))
-            : std::nullopt,
-        graph_view.number_of_vertices(),
-        k_);
-
       cugraph::unrenumber_int_vertices<vertex_t, multi_gpu>(
         handle_,
         result_src.data(),
@@ -127,9 +113,12 @@ struct k_truss_functor : public cugraph::c_api::abstract_functor {
       result_ = new cugraph::c_api::cugraph_induced_subgraph_result_t{
         new cugraph::c_api::cugraph_type_erased_device_array_t(result_src, graph_->vertex_type_),
         new cugraph::c_api::cugraph_type_erased_device_array_t(result_dst, graph_->vertex_type_),
+        /*
         wgt ? new cugraph::c_api::cugraph_type_erased_device_array_t(*result_wgt,
                                                                      graph_->weight_type_)
             : NULL,
+        */
+        NULL,
         NULL,
         NULL,
         new cugraph::c_api::cugraph_type_erased_device_array_t(edge_offsets,

cpp/src/community/edge_triangle_count_impl.cuh

@@ -0,0 +1,169 @@
+/*
+ * Copyright (c) 2020-2023, 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/detail/decompress_edge_partition.cuh>
+#include <cugraph/graph_functions.hpp>
+#include <cugraph/graph_view.hpp>
+#include <cugraph/partition_manager.hpp>
+#include <cugraph/utilities/error.hpp>
+#include <cugraph/utilities/mask_utils.cuh>
+
+#include <raft/core/handle.hpp>
+#include <raft/util/cudart_utils.hpp>
+
+#include <rmm/device_uvector.hpp>
+
+#include <thrust/adjacent_difference.h>
+#include <thrust/binary_search.h>
+#include <thrust/functional.h>
+#include <thrust/iterator/zip_iterator.h>
+#include <thrust/sort.h>
+#include <thrust/tuple.h>
+
+#include <detail/graph_partition_utils.cuh>
+#include <prims/transform_reduce_dst_nbr_intersection_of_e_endpoints_by_v.cuh>
+
+#include <optional>
+#include <tuple>
+#include <type_traits>
+#include <vector>
+
+namespace cugraph {
+
+namespace detail {
+
+template <typename vertex_t, typename edge_t, typename EdgeIterator>
+struct update_edges_p_r_q_r_num_triangles {
+  size_t num_edges{};  // rename to num_edges
+  const edge_t edge_first_or_second{};
+  raft::device_span<size_t const> intersection_offsets{};
+  raft::device_span<vertex_t const> intersection_indices{};
+  raft::device_span<edge_t> num_triangles{};
+
+  EdgeIterator edge_first{};
+
+  __device__ void operator()(size_t i) const
+  {
+    auto itr = thrust::upper_bound(
+      thrust::seq, intersection_offsets.begin() + 1, intersection_offsets.end(), i);
+    auto idx = thrust::distance(intersection_offsets.begin() + 1, itr);
+    if (edge_first_or_second == 0) {
+      auto p_r_pair =
+        thrust::make_tuple(thrust::get<0>(*(edge_first + idx)), intersection_indices[i]);
+
+      // Find its position in 'edges'
+      auto itr_p_r_p_q =
+        thrust::lower_bound(thrust::seq,
+                            edge_first,
+                            edge_first + num_edges,  // pass the number of vertex pairs
+                            p_r_pair);
+
+      assert(*itr_p_r_p_q == p_r_pair);
+      idx = thrust::distance(edge_first, itr_p_r_p_q);
+    } else {
+      auto p_r_pair =
+        thrust::make_tuple(thrust::get<1>(*(edge_first + idx)), intersection_indices[i]);
+
+      // Find its position in 'edges'
+      auto itr_p_r_p_q =
+        thrust::lower_bound(thrust::seq,
+                            edge_first,
+                            edge_first + num_edges,  // pass the number of vertex pairs
+                            p_r_pair);
+      assert(*itr_p_r_p_q == p_r_pair);
+      idx = thrust::distance(edge_first, itr_p_r_p_q);
+    }
+    cuda::atomic_ref<edge_t, cuda::thread_scope_device> atomic_counter(num_triangles[idx]);
+    auto r = atomic_counter.fetch_add(edge_t{1}, cuda::std::memory_order_relaxed);
+  }
+};
+
+template <typename vertex_t, typename edge_t, bool store_transposed, bool multi_gpu>
+std::enable_if_t<!multi_gpu, rmm::device_uvector<edge_t>> edge_triangle_count_impl(
+  raft::handle_t const& handle,
+  graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu> const& graph_view,
+  raft::device_span<vertex_t> edgelist_srcs,
+  raft::device_span<vertex_t> edgelist_dsts)
+{
+  auto edge_first = thrust::make_zip_iterator(edgelist_srcs.begin(), edgelist_dsts.begin());
+
+  thrust::sort(handle.get_thrust_policy(), edge_first, edge_first + edgelist_srcs.size());
+
+  auto [intersection_offsets, intersection_indices] =
+    detail::nbr_intersection(handle,
+                             graph_view,
+                             cugraph::edge_dummy_property_t{}.view(),
+                             edge_first,
+                             edge_first + edgelist_srcs.size(),
+                             std::array<bool, 2>{true, true},
+                             false /*FIXME: pass 'do_expensive_check' as argument*/);
+
+  // FIXME: invalid type for unsigned integers.
+  rmm::device_uvector<edge_t> num_triangles(edgelist_srcs.size(), handle.get_stream());
+
+  // Update the number of triangles of each (p, q) edges by looking at their intersection
+  // size
+  thrust::adjacent_difference(handle.get_thrust_policy(),
+                              intersection_offsets.begin() + 1,
+                              intersection_offsets.end(),
+                              num_triangles.begin());
+
+  // Given intersection offsets and indices that are used to update the number of
+  // triangles of (p, q) edges where `r`s are the intersection indices, update
+  // the number of triangles of the pairs (p, r) and (q, r).
+
+  thrust::for_each(
+    handle.get_thrust_policy(),
+    thrust::make_counting_iterator<edge_t>(0),
+    thrust::make_counting_iterator<edge_t>(intersection_indices.size()),
+    update_edges_p_r_q_r_num_triangles<vertex_t, edge_t, decltype(edge_first)>{
+      edgelist_srcs.size(),
+      0,
+      raft::device_span<size_t const>(intersection_offsets.data(), intersection_offsets.size()),
+      raft::device_span<vertex_t const>(intersection_indices.data(), intersection_indices.size()),
+      raft::device_span<edge_t>(num_triangles.data(), num_triangles.size()),
+      edge_first});
+
+  thrust::for_each(
+    handle.get_thrust_policy(),
+    thrust::make_counting_iterator<edge_t>(0),
+    thrust::make_counting_iterator<edge_t>(intersection_indices.size()),
+    update_edges_p_r_q_r_num_triangles<vertex_t, edge_t, decltype(edge_first)>{
+      edgelist_srcs.size(),
+      1,
+      raft::device_span<size_t const>(intersection_offsets.data(), intersection_offsets.size()),
+      raft::device_span<vertex_t const>(intersection_indices.data(), intersection_indices.size()),
+      raft::device_span<edge_t>(num_triangles.data(), num_triangles.size()),
+      edge_first});
+
+  return num_triangles;
+}
+
+}  // namespace detail
+
+template <typename vertex_t, typename edge_t, bool store_transposed, bool multi_gpu>
+rmm::device_uvector<edge_t> edge_triangle_count(
+  raft::handle_t const& handle,
+  graph_view_t<vertex_t, edge_t, store_transposed, multi_gpu> const& graph_view,
+  raft::device_span<vertex_t> edgelist_srcs,
+  raft::device_span<vertex_t> edgelist_dsts)
+{
+  return detail::edge_triangle_count_impl(handle, graph_view, edgelist_srcs, edgelist_dsts);
+}
+
+}  // namespace cugraph

cpp/src/community/edge_triangle_count_sg.cu

@@ -0,0 +1,59 @@
+/*
+ * Copyright (c) 2021-2023, 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 <community/edge_triangle_count_impl.cuh>
+
+namespace cugraph {
+
+// SG instantiation
+template rmm::device_uvector<int32_t> edge_triangle_count(
+  raft::handle_t const& handle,
+  cugraph::graph_view_t<int32_t, int32_t, false, false> const& graph_view,
+  raft::device_span<int32_t> edgelist_srcs,
+  raft::device_span<int32_t> edgelist_dsts);
+
+template rmm::device_uvector<int64_t> edge_triangle_count(
+  raft::handle_t const& handle,
+  cugraph::graph_view_t<int32_t, int64_t, false, false> const& graph_view,
+  raft::device_span<int32_t> edgelist_srcs,
+  raft::device_span<int32_t> edgelist_dsts);
+
+template rmm::device_uvector<int64_t> edge_triangle_count(
+  raft::handle_t const& handle,
+  cugraph::graph_view_t<int64_t, int64_t, false, false> const& graph_view,
+  raft::device_span<int64_t> edgelist_srcs,
+  raft::device_span<int64_t> edgelist_dsts);
+
+template rmm::device_uvector<int32_t> edge_triangle_count(
+  raft::handle_t const& handle,
+  cugraph::graph_view_t<int64_t, int32_t, false, false> const& graph_view,
+  raft::device_span<int64_t> edgelist_srcs,
+  raft::device_span<int64_t> edgelist_dsts);
+
+/*
+template rmm::device_uvector<int64_t> edge_triangle_count(
+  raft::handle_t const& handle,
+  cugraph::graph_view_t<int64_t, int64_t, false, false> const& graph_view,
+  raft::device_span<int64_t> edgelist_srcs,
+  raft::device_span<int64_t> edgelist_dsts);
+
+template rmm::device_uvector<int64_t> edge_triangle_count(
+  raft::handle_t const& handle,
+  cugraph::graph_view_t<int32_t, int64_t, false, false> const& graph_view,
+  raft::device_span<int64_t> edgelist_srcs,
+  raft::device_span<int64_t> edgelist_dsts);
+*/
+
+}  // namespace cugraph