[WIP] Share COO and sparse spectral helper kernels

cpp/CMakeLists.txt

@@ -1431,6 +1431,26 @@ if(NOT BUILD_CPU_ONLY)
       "${CMAKE_CURRENT_BINARY_DIR}/src/selection/select_k_inst_data_@data_abbrev@_index_@index_abbrev@.cu"
   )
 
+  generate_string_matrix(
+    compute_graph_laplacian_detail_extern_instantiations
+    MATRIX_JSON_FILE
+    "${CMAKE_CURRENT_SOURCE_DIR}/src/core/compute_graph_laplacian_detail_matrix.json"
+    ITEM_FORMAT
+    "extern template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@@semicolon@"
+    GLUE
+    "\n"
+  )
+
+  generate_string_matrix(
+    laplacian_reduce_by_key_extern_instantiations
+    MATRIX_JSON_FILE
+    "${CMAKE_CURRENT_SOURCE_DIR}/src/core/laplacian_reduce_by_key_matrix.json"
+    ITEM_FORMAT
+    "extern template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@@semicolon@"
+    GLUE
+    "\n"
+  )
+
   generate_string_matrix(
     map_detail_extern_instantiations
     MATRIX_JSON_FILE
@@ -1711,10 +1731,139 @@ if(NOT BUILD_CPU_ONLY)
       "${CMAKE_CURRENT_BINARY_DIR}/src/core/cub_mixed_scan_kernels_@instantiation_name@.cu"
   )
 
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/compute_graph_laplacian_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/compute_graph_laplacian.cu.in"
+    OUTPUT_FILE_FORMAT
+      "${CMAKE_CURRENT_BINARY_DIR}/src/core/compute_graph_laplacian_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE
+      "${CMAKE_CURRENT_SOURCE_DIR}/src/core/compute_graph_laplacian_detail_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/compute_graph_laplacian_detail.cu.in"
+    OUTPUT_FILE_FORMAT
+      "${CMAKE_CURRENT_BINARY_DIR}/src/core/compute_graph_laplacian_detail_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/laplacian_reduce_by_key_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/laplacian_reduce_by_key.cu.in"
+    OUTPUT_FILE_FORMAT
+      "${CMAKE_CURRENT_BINARY_DIR}/src/core/laplacian_reduce_by_key_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/symmetrize_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/symmetrize.cu.in"
+    OUTPUT_FILE_FORMAT "${CMAKE_CURRENT_BINARY_DIR}/src/core/symmetrize_@instantiation_name@.cu"
+  )
+
+  generate_string_matrix(
+    coo_merge_kernel_extern_instantiations
+    MATRIX_JSON_FILE
+    "${CMAKE_CURRENT_SOURCE_DIR}/src/core/coo_merge_kernel_matrix.json"
+    ITEM_FORMAT
+    "template <> __host__ __device__ auto @provider_type@::@method@() -> @pointer_type@@semicolon@"
+    GLUE
+    "\n"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/masked_matmul_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/masked_matmul.cu.in"
+    OUTPUT_FILE_FORMAT "${CMAKE_CURRENT_BINARY_DIR}/src/core/masked_matmul_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/coo_remove_scalar_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/coo_remove_scalar.cu.in"
+    OUTPUT_FILE_FORMAT
+      "${CMAKE_CURRENT_BINARY_DIR}/src/core/coo_remove_scalar_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/laplacian_normalized_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/laplacian_normalized.cu.in"
+    OUTPUT_FILE_FORMAT
+      "${CMAKE_CURRENT_BINARY_DIR}/src/core/laplacian_normalized_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/compute_duplicates_mask_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/compute_duplicates_mask.cu.in"
+    OUTPUT_FILE_FORMAT
+      "${CMAKE_CURRENT_BINARY_DIR}/src/core/compute_duplicates_mask_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE
+      "${CMAKE_CURRENT_SOURCE_DIR}/src/core/symmetrize_connectivity_graph_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/symmetrize_connectivity_graph.cu.in"
+    OUTPUT_FILE_FORMAT
+      "${CMAKE_CURRENT_BINARY_DIR}/src/core/symmetrize_connectivity_graph_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/coo_sort_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/coo_sort.cu.in"
+    OUTPUT_FILE_FORMAT "${CMAKE_CURRENT_BINARY_DIR}/src/core/coo_sort_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/coo_degree_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/coo_degree.cu.in"
+    OUTPUT_FILE_FORMAT "${CMAKE_CURRENT_BINARY_DIR}/src/core/coo_degree_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/faster_dot_on_csr_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/faster_dot_on_csr.cu.in"
+    OUTPUT_FILE_FORMAT
+      "${CMAKE_CURRENT_BINARY_DIR}/src/core/faster_dot_on_csr_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/fill_indices_by_rows_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/fill_indices_by_rows.cu.in"
+    OUTPUT_FILE_FORMAT
+      "${CMAKE_CURRENT_BINARY_DIR}/src/core/fill_indices_by_rows_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/laplacian_zero_to_one_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/laplacian_zero_to_one.cu.in"
+    OUTPUT_FILE_FORMAT
+      "${CMAKE_CURRENT_BINARY_DIR}/src/core/laplacian_zero_to_one_@instantiation_name@.cu"
+  )
+
+  generate_inst_matrix(
+    cuvs_primitive_sources
+    MATRIX_JSON_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/coo_merge_kernel_matrix.json"
+    INPUT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/src/core/coo_merge_kernel.cu.in"
+    OUTPUT_FILE_FORMAT
+      "${CMAKE_CURRENT_BINARY_DIR}/src/core/coo_merge_kernel_@instantiation_name@.cu"
+  )
+
   add_library(
     cuvs_objs OBJECT
     ${cuvs_primitive_sources}
     src/core/calc_nnz_by_rows.cu
+    src/core/coo_symmetrize.cu
     src/core/copy_distance_tile.cu
     src/core/cub_max.cu
     src/core/cub_ptx_query.cu

cpp/src/cluster/detail/single_linkage.cuh

@@ -6,6 +6,7 @@
 #pragma once
 
 #include "../../core/map_launchers.hpp"
+#include "../../core/spectral_launchers.hpp"
 #include "../../neighbors/detail/reachability.cuh"
 #include "agglomerative.cuh"
 #include "connectivities.cuh"
@@ -82,14 +83,14 @@ void build_mr_linkage(
       auto coo_rows = raft::make_device_vector<value_idx, value_idx>(handle, min_samples * m);
       cuvs::detail::map_offset(handle, coo_rows.view(), raft::div_const_op<value_idx>(min_samples));
 
-      raft::sparse::linalg::symmetrize(handle,
-                                       coo_rows.data_handle(),
-                                       inds.data_handle(),
-                                       dists.data_handle(),
-                                       static_cast<value_idx>(m),
-                                       static_cast<value_idx>(m),
-                                       static_cast<nnz_t>(min_samples * m),
-                                       mr_coo.value());
+      cuvs::detail::symmetrize(handle,
+                               coo_rows.data_handle(),
+                               inds.data_handle(),
+                               dists.data_handle(),
+                               static_cast<value_idx>(m),
+                               static_cast<value_idx>(m),
+                               static_cast<nnz_t>(min_samples * m),
+                               mr_coo.value());
     }  // scope to drop inds and dists matrices early
     auto mr_indptr = raft::make_device_vector<value_idx, value_idx>(handle, m + 1);
     raft::sparse::convert::sorted_coo_to_csr(

cpp/src/core/compute_duplicates_mask.cu.in

@@ -0,0 +1,36 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <core/cub_ptx_query.hpp>
+
+#include <core/coo_sort_launchers.hpp>
+
+#include <cuvs/core/cuda_fp16.hpp>
+#include <raft/util/cuda_utils.cuh>
+#include <raft/util/device_atomics.cuh>
+
+#include <raft/sparse/op/reduce.cuh>
+
+extern template void raft::sparse::op::detail::compute_duplicates_mask<int, size_t>(
+  int*, int const*, int const*, size_t, cudaStream_t);
+
+extern template void raft::sparse::op::detail::compute_duplicates_mask<long, size_t>(
+  long*, long const*, long const*, size_t, cudaStream_t);
+
+namespace cuvs::detail {
+
+template <typename value_idx>
+void compute_duplicates_mask(
+  value_idx* mask, const value_idx* rows, const value_idx* cols, size_t nnz, cudaStream_t stream)
+
+{
+  return raft::sparse::op::compute_duplicates_mask(mask, rows, cols, nnz, stream);
+}
+
+}  // namespace cuvs::detail
+
+// clang-format off
+template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@;
+// clang-format on

cpp/src/core/compute_duplicates_mask_matrix.json

@@ -0,0 +1,35 @@
+[
+  {
+    "return_type": "CUVS_EXPORT void",
+    "qualified_function": "cuvs::detail::compute_duplicates_mask",
+    "template_arguments": "",
+    "function_qualifiers": "",
+    "_variants": [
+      {
+        "parameter_types": "int*, int const*, int const*, unsigned long, CUstream_st*",
+        "instantiation_name": "cuvs_detail_compute_duplicates_mask_int_ptr_int_c_ptr_int_c_ptr_ulong_stream_ptr"
+      },
+      {
+        "parameter_types": "long*, long const*, long const*, unsigned long, CUstream_st*",
+        "instantiation_name": "cuvs_detail_compute_duplicates_mask_long_ptr_long_c_ptr_long_c_ptr_ulong_stream_ptr"
+      }
+    ]
+  },
+  {
+    "return_type": "void",
+    "qualified_function": "raft::sparse::op::detail::compute_duplicates_mask",
+    "function_qualifiers": "",
+    "_variants": [
+      {
+        "template_arguments": "<int, unsigned long>",
+        "parameter_types": "int*, int const*, int const*, unsigned long, cudaStream_t",
+        "instantiation_name": "raft_sparse_op_detail_compute_duplicates_mask_int_ulong"
+      },
+      {
+        "template_arguments": "<long, unsigned long>",
+        "parameter_types": "long*, long const*, long const*, unsigned long, cudaStream_t",
+        "instantiation_name": "raft_sparse_op_detail_compute_duplicates_mask_long_ulong"
+      }
+    ]
+  }
+]

cpp/src/core/compute_graph_laplacian.cu.in

@@ -0,0 +1,80 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <core/cub_ptx_query.hpp>
+
+#include <core/cub_scan_launchers.hpp>
+
+#include <core/spectral_launchers.hpp>
+
+#include <cuvs/core/cuda_fp16.hpp>
+#include <raft/util/cuda_utils.cuh>
+#include <raft/util/device_atomics.cuh>
+
+#include <raft/sparse/linalg/laplacian.cuh>
+
+extern template void raft::linalg::detail::map<false, double, int, raft::const_op<double>>(
+  cuda::stream_ref, double*, int, raft::const_op<double>);
+
+extern template void raft::linalg::detail::map<false, int, int, raft::const_op<int>>(
+  cuda::stream_ref, int*, int, raft::const_op<int>);
+
+extern template void raft::linalg::detail::map<false, float, int, raft::const_op<float>>(
+  cuda::stream_ref, float*, int, raft::const_op<float>);
+
+extern template void raft::linalg::detail::map<false, float, long, raft::const_op<float>>(
+  cuda::stream_ref, float*, long, raft::const_op<float>);
+
+extern template void raft::sparse::op::detail::coo_sort<float, int, int>(
+  bool, int, int, int, int*, int*, float*, cudaStream_t);
+
+extern template void raft::sparse::op::detail::coo_sort<float, int, long>(
+  bool, int, int, long, int*, int*, float*, cudaStream_t);
+
+extern template void raft::sparse::op::detail::coo_sort<double, int, int>(
+  bool, int, int, int, int*, int*, double*, cudaStream_t);
+
+extern template void raft::sparse::op::detail::coo_sort<double, int, long>(
+  bool, int, int, long, int*, int*, double*, cudaStream_t);
+
+extern template void raft::sparse::op::detail::coo_sort<float, int, unsigned long>(
+  bool, int, int, unsigned long, int*, int*, float*, cudaStream_t);
+
+extern template void raft::sparse::op::detail::coo_sort<float, long, unsigned long>(
+  bool, long, long, unsigned long, long*, long*, float*, cudaStream_t);
+
+namespace cuvs::detail {
+
+template <typename ElementType, typename IndptrType, typename IndicesType, typename NZType>
+raft::device_csr_matrix<std::remove_const_t<ElementType>,
+                        std::remove_const_t<IndptrType>,
+                        std::remove_const_t<IndicesType>,
+                        std::remove_const_t<NZType>>
+compute_graph_laplacian(
+  raft::resources const& res,
+  raft::device_csr_matrix_view<ElementType, IndptrType, IndicesType, NZType> input)
+
+{
+  return raft::sparse::linalg::compute_graph_laplacian(res, input);
+}
+
+template <typename ElementType, typename RowType, typename ColType, typename NZType>
+raft::device_coo_matrix<ElementType, RowType, ColType, NZType> compute_graph_laplacian(
+  raft::resources const& res,
+  raft::device_coo_matrix_view<ElementType, RowType, ColType, NZType> input)
+
+{
+  return raft::sparse::linalg::compute_graph_laplacian(res, input);
+}
+
+}  // namespace cuvs::detail
+
+// clang-format off
+@laplacian_reduce_by_key_extern_instantiations@
+
+@compute_graph_laplacian_detail_extern_instantiations@
+
+template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@;
+// clang-format on