Strongly Connected Components Part 1/2

cpp/CMakeLists.txt

@@ -297,6 +297,8 @@ set(CUGRAPH_SG_SOURCES
     src/tree/legacy/mst.cu
     src/components/weakly_connected_components_sg_v64_e64.cu
     src/components/weakly_connected_components_sg_v32_e32.cu
+    src/components/strongly_connected_components_sg_v64_e64.cu
+    src/components/strongly_connected_components_sg_v32_e32.cu
     src/components/mis_sg_v64_e64.cu
     src/components/mis_sg_v32_e32.cu
     src/components/vertex_coloring_sg_v64_e64.cu
@@ -426,6 +428,8 @@ set(CUGRAPH_MG_SOURCES
     src/centrality/betweenness_centrality_mg_v32_e32.cu
     src/components/weakly_connected_components_mg_v64_e64.cu
     src/components/weakly_connected_components_mg_v32_e32.cu
+    src/components/strongly_connected_components_mg_v64_e64.cu
+    src/components/strongly_connected_components_mg_v32_e32.cu
     src/components/mis_mg_v64_e64.cu
     src/components/mis_mg_v32_e32.cu
     src/components/vertex_coloring_mg_v64_e64.cu

cpp/include/cugraph/algorithms.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -1743,6 +1743,29 @@ void weakly_connected_components(raft::handle_t const& handle,
                                  vertex_t* components,
                                  bool do_expensive_check = false);
 
+/**
+.* @ingroup components_cpp
+ * @brief Finds (strongly-connected-)component IDs of each vertices in the input graph.
+ *
+ * Component IDs can be arbitrary integers (they can be non-consecutive and are not ordered by
+ * component size or any other criterion).
+ *
+ * @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)
+ * 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 do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
+ * @return Device vector of stronlgy connected component IDs
+ */
+template <typename vertex_t, typename edge_t, bool multi_gpu>
+rmm::device_uvector<vertex_t> strongly_connected_components(
+  raft::handle_t const& handle,
+  graph_view_t<vertex_t, edge_t, false, multi_gpu> const& graph_view,
+  bool do_expensive_check = false);
+
 /**
 .* @ingroup core_cpp
  * @brief  Identify whether the core number computation should be based off incoming edges,

cpp/include/cugraph/utilities/device_functors.cuh

@@ -1,11 +1,12 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
 
 #include <cugraph/utilities/packed_bool_utils.hpp>
 
+#include <thrust/binary_search.h>
 #include <thrust/iterator/iterator_traits.h>
 
 #include <cstddef>
@@ -170,6 +171,18 @@ struct divider_t {
   __device__ T operator()(T input) const { return input / divisor; }
 };
 
+template <typename T>
+struct segment_id_t {
+  raft::device_span<T const> segment_lasts{};
+
+  __device__ T operator()(T i) const
+  {
+    return static_cast<T>(cuda::std::distance(
+      segment_lasts.begin(),
+      thrust::upper_bound(thrust::seq, segment_lasts.begin(), segment_lasts.end(), i)));
+  }
+};
+
 }  // namespace detail
 
 }  // namespace cugraph