[FEA] Build Single Linkage API

cpp/include/cuvs/cluster/agglomerative.hpp

@@ -119,7 +119,4 @@ void single_linkage(
   cuvs::cluster::agglomerative::Linkage linkage = cuvs::cluster::agglomerative::Linkage::KNN_GRAPH,
   std::optional<int> c                          = std::make_optional<int>(DEFAULT_CONST_C));
 
-/**
- * @}
- */
 };  // end namespace  cuvs::cluster::agglomerative

cpp/include/cuvs/neighbors/reachability.hpp

@@ -16,6 +16,7 @@
 
 #pragma once
 
+#include <raft/core/device_coo_matrix.hpp>
 #include <raft/core/device_mdspan.hpp>
 #include <raft/core/handle.hpp>
 #include <raft/sparse/coo.hpp>
@@ -73,6 +74,35 @@ void mutual_reachability_graph(
   raft::sparse::COO<float, int>& out,
   cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2SqrtExpanded,
   float alpha                         = 1.0);
+
+namespace helpers {
+/**
+ * Given a mutual reachability graph, connects graph components and build dendrogram.
+ * Returns mst edges sorted by weight and the linkage.
+ * @tparam value_idx
+ * @tparam value_t
+ * @param[in] handle raft handle for resource reuse
+ * @param[in] X data points (size m * n)
+ * @param[in] metric distance metric to use
+ * @param[in] graph_indptr CSR indices of graph nodes (size m + 1)
+ * @param[in] graph input graph
+ * @param[out] out_mst output MST sorted by edge weights (size m - 1)
+ * @param[out] dendrogram output dendrogram (size [n_rows - 1] * 2)
+ * @param[out] out_distances distances for output
+ * @param[out] out_sizes cluster sizes of output
+ */
+void build_single_linkage_dendrogram(raft::resources const& handle,
+                                     raft::device_matrix_view<const float, int, raft::row_major> X,
+                                     cuvs::distance::DistanceType metric,
+                                     raft::device_vector_view<int, int> graph_indptr,
+                                     raft::device_coo_matrix_view<float, int, int, size_t> graph,
+                                     raft::device_vector_view<float, int> core_dists,
+                                     raft::device_coo_matrix_view<float, int, int, int> out_mst,
+                                     raft::device_matrix_view<int, int> dendrogram,
+                                     raft::device_vector_view<float, int> out_distances,
+                                     raft::device_vector_view<int, int> out_sizes);
+}  // namespace helpers
+
 /**
  * @}
  */

cpp/src/cluster/detail/single_linkage.cuh

@@ -122,4 +122,5 @@ void single_linkage(raft::resources const& handle,
   out->n_leaves               = m;
   out->n_connected_components = 1;
 }
+
 };  // namespace  cuvs::cluster::agglomerative::detail

cpp/src/cluster/single_linkage_float.cu

@@ -39,4 +39,5 @@ void single_linkage(raft::resources const& handle,
       handle, X, dendrogram, labels, metric, n_clusters, c);
   }
 }
+
 }  // namespace cuvs::cluster::agglomerative

cpp/src/neighbors/detail/reachability.cuh

@@ -15,6 +15,9 @@
  */
 
 #pragma once
+#include "../../cluster/detail/agglomerative.cuh"
+#include "../../cluster/detail/mst.cuh"
+#include "../../sparse/neighbors/cross_component_nn.cuh"
 #include "./knn_brute_force.cuh"
 
 #include <raft/linalg/unary_op.cuh>
@@ -254,4 +257,77 @@ void mutual_reachability_graph(const raft::resources& handle,
                     });
 }
 
+/**
+ * Given a mutual reachability graph and core distances, constructs a linkage over it by computing
+ * the minimum spanning tree and dendrogram. Returns mst edges sorted by weight and the linkage.
+ * @tparam value_idx
+ * @tparam value_t
+ * @param[in] handle raft handle for resource reuse
+ * @param[in] X data points (size m * n)
+ * @param[in] m number of rows
+ * @param[in] n number of columns
+ * @param[in] metric distance metric to use
+ * @param[in] indptr CSR indices of mutual reachability knn graph (size m + 1)
+ * @param[out] graph_coo input graph
+ * @param[out] out_mst_src src vertex of MST edges (size m - 1)
+ * @param[out] out_mst_dst dst vertex of MST eges (size m - 1)
+ * @param[out] out_mst_weights weights of MST edges (size m - 1)
+ * @param[out] out_children output dendrogram
+ * @param[out] out_deltas distances of output
+ * @param[out] out_sizes cluster sizes of output
+ */
+template <typename value_idx = int, typename value_t = float, typename nnz_t>
+void build_single_linkage_dendrogram(
+  raft::resources const& handle,
+  const value_t* X,
+  size_t m,
+  size_t n,
+  cuvs::distance::DistanceType metric,
+  value_idx* indptr,
+  raft::device_coo_matrix_view<value_t, value_idx, value_idx, nnz_t> graph_coo,
+  value_t* core_dists,
+  value_idx* out_mst_src,
+  value_idx* out_mst_dst,
+  value_t* out_mst_weights,
+  value_idx* out_dendrogram,
+  value_t* out_deltas,
+  value_idx* out_sizes)
+{
+  /**
+   * Construct MST sorted by weights
+   */
+  auto color = raft::make_device_vector<value_idx, value_idx>(handle, static_cast<value_idx>(m));
+  cuvs::sparse::neighbors::MutualReachabilityFixConnectivitiesRedOp<value_idx, value_t>
+    reduction_op(core_dists, m);
+
+  // during knn graph connection
+  cuvs::cluster::agglomerative::detail::build_sorted_mst(
+    handle,
+    X,
+    indptr,
+    graph_coo.structure_view().get_cols().data(),
+    graph_coo.get_elements().data(),
+    m,
+    n,
+    out_mst_src,
+    out_mst_dst,
+    out_mst_weights,
+    color.data_handle(),
+    graph_coo.structure_view().get_nnz(),
+    reduction_op,
+    metric);
+
+  /**
+   * Perform hierarchical labeling
+   */
+  cuvs::cluster::agglomerative::detail::build_dendrogram_host(handle,
+                                                              out_mst_src,
+                                                              out_mst_dst,
+                                                              out_mst_weights,
+                                                              m - 1,
+                                                              out_dendrogram,
+                                                              out_deltas,
+                                                              out_sizes);
+}
+
 }  // namespace cuvs::neighbors::detail::reachability

cpp/src/neighbors/reachability.cu

@@ -46,4 +46,34 @@ void mutual_reachability_graph(const raft::resources& handle,
     core_dists.data_handle(),
     out);
 }
+
+namespace helpers {
+void build_single_linkage_dendrogram(raft::resources const& handle,
+                                     raft::device_matrix_view<const float, int, raft::row_major> X,
+                                     cuvs::distance::DistanceType metric,
+                                     raft::device_vector_view<int, int> graph_indptr,
+                                     raft::device_coo_matrix_view<float, int, int, size_t> graph,
+                                     raft::device_vector_view<float, int> core_dists,
+                                     raft::device_coo_matrix_view<float, int, int, int> out_mst,
+                                     raft::device_matrix_view<int, int> dendrogram,
+                                     raft::device_vector_view<float, int> out_distances,
+                                     raft::device_vector_view<int, int> out_sizes)
+{
+  cuvs::neighbors::detail::reachability::build_single_linkage_dendrogram(
+    handle,
+    X.data_handle(),
+    static_cast<size_t>(X.extent(0)),
+    static_cast<size_t>(X.extent(1)),
+    metric,
+    graph_indptr.data_handle(),
+    graph,
+    core_dists.data_handle(),
+    out_mst.structure_view().get_rows().data(),
+    out_mst.structure_view().get_cols().data(),
+    out_mst.get_elements().data(),
+    dendrogram.data_handle(),
+    out_distances.data_handle(),
+    out_sizes.data_handle());
+}
+}  // namespace helpers
 }  // namespace cuvs::neighbors::reachability

cpp/src/sparse/neighbors/cross_component_nn.cuh

@@ -26,6 +26,10 @@ namespace cuvs::sparse::neighbors {
 template <typename value_idx, typename value_t>
 using FixConnectivitiesRedOp = detail::FixConnectivitiesRedOp<value_idx, value_t>;
 
+template <typename value_idx, typename value_t>
+using MutualReachabilityFixConnectivitiesRedOp =
+  detail::MutualReachabilityFixConnectivitiesRedOp<value_idx, value_t>;
+
 /**
  * Gets the number of unique components from array of
  * colors or labels. This does not assume the components are

cpp/src/sparse/neighbors/detail/cross_component_nn.cuh

@@ -108,6 +108,93 @@ struct FixConnectivitiesRedOp {
   void scatter(const raft::resources& handle, value_idx* map) {}
 };
 
+template <typename value_idx, typename value_t>
+struct MutualReachabilityFixConnectivitiesRedOp {
+  value_t* core_dists;
+  value_idx m;
+
+  DI MutualReachabilityFixConnectivitiesRedOp() : m(0) {}
+
+  MutualReachabilityFixConnectivitiesRedOp(value_t* core_dists_, value_idx m_)
+    : core_dists(core_dists_), m(m_){};
+
+  typedef typename raft::KeyValuePair<value_idx, value_t> KVP;
+  DI void operator()(value_idx rit, KVP* out, const KVP& other) const
+  {
+    if (rit < m && other.value < std::numeric_limits<value_t>::max()) {
+      value_t core_dist_rit   = core_dists[rit];
+      value_t core_dist_other = max(core_dist_rit, max(core_dists[other.key], other.value));
+
+      value_t core_dist_out;
+      if (out->key > -1) {
+        core_dist_out = max(core_dist_rit, max(core_dists[out->key], out->value));
+      } else {
+        core_dist_out = out->value;
+      }
+
+      bool smaller = core_dist_other < core_dist_out;
+      out->key     = smaller ? other.key : out->key;
+      out->value   = smaller ? core_dist_other : core_dist_out;
+    }
+  }
+
+  DI KVP operator()(value_idx rit, const KVP& a, const KVP& b) const
+  {
+    if (rit < m && a.key > -1) {
+      value_t core_dist_rit = core_dists[rit];
+      value_t core_dist_a   = max(core_dist_rit, max(core_dists[a.key], a.value));
+
+      value_t core_dist_b;
+      if (b.key > -1) {
+        core_dist_b = max(core_dist_rit, max(core_dists[b.key], b.value));
+      } else {
+        core_dist_b = b.value;
+      }
+
+      return core_dist_a < core_dist_b ? KVP(a.key, core_dist_a) : KVP(b.key, core_dist_b);
+    }
+
+    return b;
+  }
+
+  DI void init(value_t* out, value_t maxVal) const { *out = maxVal; }
+  DI void init(KVP* out, value_t maxVal) const
+  {
+    out->key   = -1;
+    out->value = maxVal;
+  }
+
+  DI void init_key(value_t& out, value_idx idx) const { return; }
+  DI void init_key(KVP& out, value_idx idx) const { out.key = idx; }
+
+  DI value_t get_value(KVP& out) const { return out.value; }
+  DI value_t get_value(value_t& out) const { return out; }
+
+  void gather(const raft::resources& handle, value_idx* map)
+  {
+    auto tmp_core_dists = raft::make_device_vector<value_t>(handle, m);
+    thrust::gather(raft::resource::get_thrust_policy(handle),
+                   map,
+                   map + m,
+                   core_dists,
+                   tmp_core_dists.data_handle());
+    raft::copy_async(
+      core_dists, tmp_core_dists.data_handle(), m, raft::resource::get_cuda_stream(handle));
+  }
+
+  void scatter(const raft::resources& handle, value_idx* map)
+  {
+    auto tmp_core_dists = raft::make_device_vector<value_t>(handle, m);
+    thrust::scatter(raft::resource::get_thrust_policy(handle),
+                    core_dists,
+                    core_dists + m,
+                    map,
+                    tmp_core_dists.data_handle());
+    raft::copy_async(
+      core_dists, tmp_core_dists.data_handle(), m, raft::resource::get_cuda_stream(handle));
+  }
+};
+
 /**
  * Assumes 3-iterator tuple containing COO rows, cols, and
  * a cub keyvalue pair object. Sorts the 3 arrays in

cpp/tests/sparse/neighbors/cross_component_nn.cu

@@ -384,93 +384,6 @@ INSTANTIATE_TEST_CASE_P(ConnectComponentsTest,
                         ConnectComponentsTestF_Int,
                         ::testing::ValuesIn(fix_conn_inputsf2));
 
-template <typename value_idx, typename value_t>
-struct MutualReachabilityFixConnectivitiesRedOp {
-  value_t* core_dists;
-  value_idx m;
-
-  DI MutualReachabilityFixConnectivitiesRedOp() : m(0) {}
-
-  MutualReachabilityFixConnectivitiesRedOp(value_t* core_dists_, value_idx m_)
-    : core_dists(core_dists_), m(m_){};
-
-  typedef typename raft::KeyValuePair<value_idx, value_t> KVP;
-  DI void operator()(value_idx rit, KVP* out, const KVP& other) const
-  {
-    if (rit < m && other.value < std::numeric_limits<value_t>::max()) {
-      value_t core_dist_rit   = core_dists[rit];
-      value_t core_dist_other = max(core_dist_rit, max(core_dists[other.key], other.value));
-
-      value_t core_dist_out;
-      if (out->key > -1) {
-        core_dist_out = max(core_dist_rit, max(core_dists[out->key], out->value));
-      } else {
-        core_dist_out = out->value;
-      }
-
-      bool smaller = core_dist_other < core_dist_out;
-      out->key     = smaller ? other.key : out->key;
-      out->value   = smaller ? core_dist_other : core_dist_out;
-    }
-  }
-
-  DI KVP operator()(value_idx rit, const KVP& a, const KVP& b) const
-  {
-    if (rit < m && a.key > -1) {
-      value_t core_dist_rit = core_dists[rit];
-      value_t core_dist_a   = max(core_dist_rit, max(core_dists[a.key], a.value));
-
-      value_t core_dist_b;
-      if (b.key > -1) {
-        core_dist_b = max(core_dist_rit, max(core_dists[b.key], b.value));
-      } else {
-        core_dist_b = b.value;
-      }
-
-      return core_dist_a < core_dist_b ? KVP(a.key, core_dist_a) : KVP(b.key, core_dist_b);
-    }
-
-    return b;
-  }
-
-  DI void init(value_t* out, value_t maxVal) const { *out = maxVal; }
-  DI void init(KVP* out, value_t maxVal) const
-  {
-    out->key   = -1;
-    out->value = maxVal;
-  }
-
-  DI void init_key(value_t& out, value_idx idx) const { return; }
-  DI void init_key(KVP& out, value_idx idx) const { out.key = idx; }
-
-  DI value_t get_value(KVP& out) const { return out.value; }
-  DI value_t get_value(value_t& out) const { return out; }
-
-  void gather(const raft::resources& handle, value_idx* map)
-  {
-    auto tmp_core_dists = raft::make_device_vector<value_t>(handle, m);
-    thrust::gather(raft::resource::get_thrust_policy(handle),
-                   map,
-                   map + m,
-                   core_dists,
-                   tmp_core_dists.data_handle());
-    raft::copy_async(
-      core_dists, tmp_core_dists.data_handle(), m, raft::resource::get_cuda_stream(handle));
-  }
-
-  void scatter(const raft::resources& handle, value_idx* map)
-  {
-    auto tmp_core_dists = raft::make_device_vector<value_t>(handle, m);
-    thrust::scatter(raft::resource::get_thrust_policy(handle),
-                    core_dists,
-                    core_dists + m,
-                    map,
-                    tmp_core_dists.data_handle());
-    raft::copy_async(
-      core_dists, tmp_core_dists.data_handle(), m, raft::resource::get_cuda_stream(handle));
-  }
-};
-
 template <typename value_t, typename value_idx>
 struct ConnectComponentsMutualReachabilityInputs {
   value_idx n_row;
@@ -518,8 +431,8 @@ class ConnectComponentsEdgesTest
     /**
      * 3. cross_component_nn to fix connectivities
      */
-    MutualReachabilityFixConnectivitiesRedOp<value_idx, value_t> red_op(core_dists.data(),
-                                                                        params.n_row);
+    cuvs::sparse::neighbors::MutualReachabilityFixConnectivitiesRedOp<value_idx, value_t> red_op(
+      core_dists.data(), params.n_row);
 
     cuvs::sparse::neighbors::cross_component_nn<value_idx, value_t>(handle,
                                                                     out_edges_unbatched,