diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index aa17055135..30dd5658a4 100644 --- a/cpp/CMakeLists.txt +++ b/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 diff --git a/cpp/src/cluster/detail/single_linkage.cuh b/cpp/src/cluster/detail/single_linkage.cuh index 83ba6cdc3f..c02d87e0ec 100644 --- a/cpp/src/cluster/detail/single_linkage.cuh +++ b/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(handle, min_samples * m); cuvs::detail::map_offset(handle, coo_rows.view(), raft::div_const_op(min_samples)); - raft::sparse::linalg::symmetrize(handle, - coo_rows.data_handle(), - inds.data_handle(), - dists.data_handle(), - static_cast(m), - static_cast(m), - static_cast(min_samples * m), - mr_coo.value()); + cuvs::detail::symmetrize(handle, + coo_rows.data_handle(), + inds.data_handle(), + dists.data_handle(), + static_cast(m), + static_cast(m), + static_cast(min_samples * m), + mr_coo.value()); } // scope to drop inds and dists matrices early auto mr_indptr = raft::make_device_vector(handle, m + 1); raft::sparse::convert::sorted_coo_to_csr( diff --git a/cpp/src/core/compute_duplicates_mask.cu.in b/cpp/src/core/compute_duplicates_mask.cu.in new file mode 100644 index 0000000000..c850ad06c4 --- /dev/null +++ b/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 + +#include + +#include +#include +#include + +#include + +extern template void raft::sparse::op::detail::compute_duplicates_mask( + int*, int const*, int const*, size_t, cudaStream_t); + +extern template void raft::sparse::op::detail::compute_duplicates_mask( + long*, long const*, long const*, size_t, cudaStream_t); + +namespace cuvs::detail { + +template +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 diff --git a/cpp/src/core/compute_duplicates_mask_matrix.json b/cpp/src/core/compute_duplicates_mask_matrix.json new file mode 100644 index 0000000000..1097c246ee --- /dev/null +++ b/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": "", + "parameter_types": "int*, int const*, int const*, unsigned long, cudaStream_t", + "instantiation_name": "raft_sparse_op_detail_compute_duplicates_mask_int_ulong" + }, + { + "template_arguments": "", + "parameter_types": "long*, long const*, long const*, unsigned long, cudaStream_t", + "instantiation_name": "raft_sparse_op_detail_compute_duplicates_mask_long_ulong" + } + ] + } +] diff --git a/cpp/src/core/compute_graph_laplacian.cu.in b/cpp/src/core/compute_graph_laplacian.cu.in new file mode 100644 index 0000000000..d0949c8c5a --- /dev/null +++ b/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 + +#include + +#include + +#include +#include +#include + +#include + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, double*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, int*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, float*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, float*, long, raft::const_op); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, int, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, long, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, int, int*, int*, double*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, long, int*, int*, double*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, unsigned long, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, long, long, unsigned long, long*, long*, float*, cudaStream_t); + +namespace cuvs::detail { + +template +raft::device_csr_matrix, + std::remove_const_t, + std::remove_const_t, + std::remove_const_t> +compute_graph_laplacian( + raft::resources const& res, + raft::device_csr_matrix_view input) + +{ + return raft::sparse::linalg::compute_graph_laplacian(res, input); +} + +template +raft::device_coo_matrix compute_graph_laplacian( + raft::resources const& res, + raft::device_coo_matrix_view 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 diff --git a/cpp/src/core/compute_graph_laplacian_detail.cu.in b/cpp/src/core/compute_graph_laplacian_detail.cu.in new file mode 100644 index 0000000000..21c287082c --- /dev/null +++ b/cpp/src/core/compute_graph_laplacian_detail.cu.in @@ -0,0 +1,52 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include + +#include + +#include +#include +#include + +#include + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, double*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, int*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, float*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, float*, long, raft::const_op); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, int, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, long, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, int, int*, int*, double*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, long, int*, int*, double*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, unsigned long, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, long, long, unsigned long, long*, long*, float*, cudaStream_t); + +// clang-format off +@laplacian_reduce_by_key_extern_instantiations@ + +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on diff --git a/cpp/src/core/compute_graph_laplacian_detail_matrix.json b/cpp/src/core/compute_graph_laplacian_detail_matrix.json new file mode 100644 index 0000000000..b1924c96a7 --- /dev/null +++ b/cpp/src/core/compute_graph_laplacian_detail_matrix.json @@ -0,0 +1,34 @@ +[ + { + "return_type": "raft::device_coo_matrix", + "qualified_function": "raft::sparse::linalg::detail::compute_graph_laplacian", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::device_coo_matrix_view", + "function_qualifiers": "", + "instantiation_name": "raft_sparse_linalg_detail_compute_graph_laplacian_f32_int_int_int" + }, + { + "return_type": "raft::device_coo_matrix", + "qualified_function": "raft::sparse::linalg::detail::compute_graph_laplacian", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::device_coo_matrix_view", + "function_qualifiers": "", + "instantiation_name": "raft_sparse_linalg_detail_compute_graph_laplacian_f32_int_int_long" + }, + { + "return_type": "raft::device_coo_matrix", + "qualified_function": "raft::sparse::linalg::detail::compute_graph_laplacian", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::device_coo_matrix_view", + "function_qualifiers": "", + "instantiation_name": "raft_sparse_linalg_detail_compute_graph_laplacian_f64_int_int_int" + }, + { + "return_type": "raft::device_coo_matrix", + "qualified_function": "raft::sparse::linalg::detail::compute_graph_laplacian", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::device_coo_matrix_view", + "function_qualifiers": "", + "instantiation_name": "raft_sparse_linalg_detail_compute_graph_laplacian_f64_int_int_long" + } +] diff --git a/cpp/src/core/compute_graph_laplacian_matrix.json b/cpp/src/core/compute_graph_laplacian_matrix.json new file mode 100644 index 0000000000..364a20dae8 --- /dev/null +++ b/cpp/src/core/compute_graph_laplacian_matrix.json @@ -0,0 +1,34 @@ +[ + { + "return_type": "raft::coo_matrix< double, int, int, int, true, raft::device_container_policy, (raft::SparsityType)0, std::conditional, raft::coordinate_structure_view>::type>", + "qualified_function": "cuvs::detail::compute_graph_laplacian", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::coo_matrix_view", + "function_qualifiers": "", + "instantiation_name": "cuvs_detail_compute_graph_laplacian_coo_matrix_view_f64_int_int_int_true" + }, + { + "return_type": "raft::coo_matrix< double, int, int, long, true, raft::device_container_policy, (raft::SparsityType)0, std::conditional, raft::coordinate_structure_view>::type>", + "qualified_function": "cuvs::detail::compute_graph_laplacian", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::coo_matrix_view", + "function_qualifiers": "", + "instantiation_name": "cuvs_detail_compute_graph_laplacian_coo_matrix_view_f64_int_int_long_true" + }, + { + "return_type": "raft::coo_matrix< float, int, int, int, true, raft::device_container_policy, (raft::SparsityType)0, std::conditional, raft::coordinate_structure_view>::type>", + "qualified_function": "cuvs::detail::compute_graph_laplacian", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::coo_matrix_view", + "function_qualifiers": "", + "instantiation_name": "cuvs_detail_compute_graph_laplacian_coo_matrix_view_f32_int_int_int_true" + }, + { + "return_type": "raft::coo_matrix< float, int, int, long, true, raft::device_container_policy, (raft::SparsityType)0, std::conditional, raft::coordinate_structure_view>::type>", + "qualified_function": "cuvs::detail::compute_graph_laplacian", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::coo_matrix_view", + "function_qualifiers": "", + "instantiation_name": "cuvs_detail_compute_graph_laplacian_coo_matrix_view_f32_int_int_long_true" + } +] diff --git a/cpp/src/core/coo_degree.cu.in b/cpp/src/core/coo_degree.cu.in new file mode 100644 index 0000000000..6111ddcde2 --- /dev/null +++ b/cpp/src/core/coo_degree.cu.in @@ -0,0 +1,14 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include + +#include + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on diff --git a/cpp/src/core/coo_degree_matrix.json b/cpp/src/core/coo_degree_matrix.json new file mode 100644 index 0000000000..9afc2d8f58 --- /dev/null +++ b/cpp/src/core/coo_degree_matrix.json @@ -0,0 +1,19 @@ +[ + { + "return_type": "void", + "qualified_function": "raft::sparse::linalg::detail::coo_degree", + "function_qualifiers": "", + "_variants": [ + { + "template_arguments": "<64, int, int, int>", + "parameter_types": "int const*, int, int*, cudaStream_t", + "instantiation_name": "64_int_int_int" + }, + { + "template_arguments": "<64, int, long, long>", + "parameter_types": "int const*, long, long*, cudaStream_t", + "instantiation_name": "64_int_long_long" + } + ] + } +] diff --git a/cpp/src/core/coo_merge_kernel.cu.in b/cpp/src/core/coo_merge_kernel.cu.in new file mode 100644 index 0000000000..0a0e5c4ebe --- /dev/null +++ b/cpp/src/core/coo_merge_kernel.cu.in @@ -0,0 +1,24 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include + +#include + +namespace { +using provider_t = @provider_type@; +using pointer_t = @pointer_type@; +} // namespace + +// clang-format off +@coo_merge_kernel_extern_instantiations@ +template <> +__host__ __device__ pointer_t provider_t::@method@() +{ + return @kernel_expression@; +} +// clang-format on diff --git a/cpp/src/core/coo_merge_kernel_matrix.json b/cpp/src/core/coo_merge_kernel_matrix.json new file mode 100644 index 0000000000..647b1b08df --- /dev/null +++ b/cpp/src/core/coo_merge_kernel_matrix.json @@ -0,0 +1,86 @@ +[ + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned long long, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, float*, unsigned long long, cuda::std::tuple*, float*, raft::sparse::op::detail::TupleComp, unsigned long long*, unsigned long long, cub::detail::vsmem_t)", + "method": "MergeSortMergeKernel", + "kernel_expression": "&cub::detail::merge_sort::DeviceMergeSortMergeKernel>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned long long, raft::sparse::op::detail::TupleComp, cuda::std::tuple, float>", + "instantiation_name": "i32_f32_u64_MergeSortMergeKernel" + }, + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned long long, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, cuda::std::tuple*, unsigned long long, unsigned long long, unsigned long long*, raft::sparse::op::detail::TupleComp, unsigned long long, int)", + "method": "MergeSortPartitionKernel", + "kernel_expression": "&cub::detail::merge_sort::DeviceMergeSortPartitionKernel>, unsigned long long, raft::sparse::op::detail::TupleComp, cuda::std::tuple>", + "instantiation_name": "i32_f32_u64_MergeSortPartitionKernel" + }, + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned long long, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned long long, cuda::std::tuple*, float*, raft::sparse::op::detail::TupleComp, cub::detail::vsmem_t)", + "method": "MergeSortBlockSortKernel", + "kernel_expression": "&cub::detail::merge_sort::DeviceMergeSortBlockSortKernel>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned long long, raft::sparse::op::detail::TupleComp, cuda::std::tuple, float>", + "instantiation_name": "i32_f32_u64_MergeSortBlockSortKernel" + }, + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned int, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, float*, unsigned int, cuda::std::tuple*, float*, raft::sparse::op::detail::TupleComp, unsigned int*, unsigned int, cub::detail::vsmem_t)", + "method": "MergeSortMergeKernel", + "kernel_expression": "&cub::detail::merge_sort::DeviceMergeSortMergeKernel>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned int, raft::sparse::op::detail::TupleComp, cuda::std::tuple, float>", + "instantiation_name": "i32_f32_u32_MergeSortMergeKernel" + }, + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned int, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, cuda::std::tuple*, unsigned int, unsigned int, unsigned int*, raft::sparse::op::detail::TupleComp, unsigned int, int)", + "method": "MergeSortPartitionKernel", + "kernel_expression": "&cub::detail::merge_sort::DeviceMergeSortPartitionKernel>, unsigned int, raft::sparse::op::detail::TupleComp, cuda::std::tuple>", + "instantiation_name": "i32_f32_u32_MergeSortPartitionKernel" + }, + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned int, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned int, cuda::std::tuple*, float*, raft::sparse::op::detail::TupleComp, cub::detail::vsmem_t)", + "method": "MergeSortBlockSortKernel", + "kernel_expression": "&cub::detail::merge_sort::DeviceMergeSortBlockSortKernel>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned int, raft::sparse::op::detail::TupleComp, cuda::std::tuple, float>", + "instantiation_name": "i32_f32_u32_MergeSortBlockSortKernel" + }, + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned long long, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, double*, unsigned long long, cuda::std::tuple*, double*, raft::sparse::op::detail::TupleComp, unsigned long long*, unsigned long long, cub::detail::vsmem_t)", + "method": "MergeSortMergeKernel", + "kernel_expression": "&cub::detail::merge_sort::DeviceMergeSortMergeKernel>>, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned long long, raft::sparse::op::detail::TupleComp, cuda::std::tuple, double>", + "instantiation_name": "i32_f64_u64_MergeSortMergeKernel" + }, + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned long long, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, cuda::std::tuple*, unsigned long long, unsigned long long, unsigned long long*, raft::sparse::op::detail::TupleComp, unsigned long long, int)", + "method": "MergeSortPartitionKernel", + "kernel_expression": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned long long, raft::sparse::op::detail::TupleComp>::MergeSortPartitionKernel()", + "instantiation_name": "i32_f64_u64_MergeSortPartitionKernel" + }, + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned long long, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned long long, cuda::std::tuple*, double*, raft::sparse::op::detail::TupleComp, cub::detail::vsmem_t)", + "method": "MergeSortBlockSortKernel", + "kernel_expression": "&cub::detail::merge_sort::DeviceMergeSortBlockSortKernel>>, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned long long, raft::sparse::op::detail::TupleComp, cuda::std::tuple, double>", + "instantiation_name": "i32_f64_u64_MergeSortBlockSortKernel" + }, + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned int, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, double*, unsigned int, cuda::std::tuple*, double*, raft::sparse::op::detail::TupleComp, unsigned int*, unsigned int, cub::detail::vsmem_t)", + "method": "MergeSortMergeKernel", + "kernel_expression": "&cub::detail::merge_sort::DeviceMergeSortMergeKernel>>, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned int, raft::sparse::op::detail::TupleComp, cuda::std::tuple, double>", + "instantiation_name": "i32_f64_u32_MergeSortMergeKernel" + }, + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned int, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, cuda::std::tuple*, unsigned int, unsigned int, unsigned int*, raft::sparse::op::detail::TupleComp, unsigned int, int)", + "method": "MergeSortPartitionKernel", + "kernel_expression": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, float*, thrust::zip_iterator>, float*, unsigned int, raft::sparse::op::detail::TupleComp>::MergeSortPartitionKernel()", + "instantiation_name": "i32_f64_u32_MergeSortPartitionKernel" + }, + { + "provider_type": "cub::detail::merge_sort::DeviceMergeSortKernelSource>>, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned int, raft::sparse::op::detail::TupleComp>", + "pointer_type": "void (*)(bool, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned int, cuda::std::tuple*, double*, raft::sparse::op::detail::TupleComp, cub::detail::vsmem_t)", + "method": "MergeSortBlockSortKernel", + "kernel_expression": "&cub::detail::merge_sort::DeviceMergeSortBlockSortKernel>>, thrust::zip_iterator>, double*, thrust::zip_iterator>, double*, unsigned int, raft::sparse::op::detail::TupleComp, cuda::std::tuple, double>", + "instantiation_name": "i32_f64_u32_MergeSortBlockSortKernel" + } +] diff --git a/cpp/src/core/coo_remove_scalar.cu.in b/cpp/src/core/coo_remove_scalar.cu.in new file mode 100644 index 0000000000..ff03e88382 --- /dev/null +++ b/cpp/src/core/coo_remove_scalar.cu.in @@ -0,0 +1,99 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include + +#include + +#include + +#include +#include +#include + +#include + +extern template thrust::device_ptr thrust::cuda_cub::exclusive_scan, + thrust::device_ptr, + int, + ::cuda::std::plus>( + thrust::cuda_cub::execution_policy&, + thrust::device_ptr, + thrust::device_ptr, + thrust::device_ptr, + int, + ::cuda::std::plus); + +extern template thrust::device_ptr thrust::cuda_cub::exclusive_scan, + thrust::device_ptr, + long, + ::cuda::std::plus>( + thrust::cuda_cub::execution_policy&, + thrust::device_ptr, + thrust::device_ptr, + thrust::device_ptr, + long, + ::cuda::std::plus); + +extern template void raft::sparse::linalg::detail::coo_degree<64, int, int, int>(int const*, + int, + int*, + cudaStream_t); + +extern template void raft::sparse::linalg::detail::coo_degree<64, int, long, long>(int const*, + long, + long*, + cudaStream_t); + +namespace cuvs::detail { + +template +void coo_remove_scalar(const idx_t* rows, + const idx_t* cols, + const T* vals, + nnz_t nnz, + idx_t* crows, + idx_t* ccols, + T* cvals, + nnz_t* cnnz, + nnz_t* cur_cnnz, + T scalar, + idx_t n, + cudaStream_t stream) + +{ + return raft::sparse::op::coo_remove_scalar( + rows, cols, vals, nnz, crows, ccols, cvals, cnnz, cur_cnnz, scalar, n, stream); +} + +template +void coo_remove_scalar(raft::sparse::COO* in, + raft::sparse::COO* out, + T scalar, + cudaStream_t stream) + +{ + return raft::sparse::op::coo_remove_scalar(in, out, scalar, stream); +} + +template +void coo_remove_scalar(raft::resources const& handle, + raft::device_coo_matrix_view in, + raft::host_scalar_view scalar, + raft::device_coo_matrix& out) + +{ + return raft::sparse::op::coo_remove_scalar(handle, in, scalar, out); +} + +} // namespace cuvs::detail + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on diff --git a/cpp/src/core/coo_remove_scalar_matrix.json b/cpp/src/core/coo_remove_scalar_matrix.json new file mode 100644 index 0000000000..7b1ac24a68 --- /dev/null +++ b/cpp/src/core/coo_remove_scalar_matrix.json @@ -0,0 +1,18 @@ +[ + { + "return_type": "void", + "qualified_function": "cuvs::detail::coo_remove_scalar", + "template_arguments": "<128>", + "function_qualifiers": "", + "_variants": [ + { + "parameter_types": "raft::resources const&, raft::coo_matrix_view, cuda::std::mdspan< float const, cuda::std::extents, cuda::std::layout_right, raft::host_device_accessor, (raft::memory_type)0>>, raft::coo_matrix< float, int, int, int, true, raft::device_container_policy, (raft::SparsityType)0, std::conditional, raft::coordinate_structure_view>::type>&", + "instantiation_name": "tpb128_f32_i32_nnz_i32" + }, + { + "parameter_types": "raft::resources const&, raft::coo_matrix_view, cuda::std::mdspan< float const, cuda::std::extents, cuda::std::layout_right, raft::host_device_accessor, (raft::memory_type)0>>, raft::coo_matrix< float, int, int, long, true, raft::device_container_policy, (raft::SparsityType)0, std::conditional< true, raft::coordinate_structure, raft::coordinate_structure_view>::type>&", + "instantiation_name": "tpb128_f32_i32_nnz_i64" + } + ] + } +] diff --git a/cpp/src/core/coo_sort.cu.in b/cpp/src/core/coo_sort.cu.in new file mode 100644 index 0000000000..7849e03b2a --- /dev/null +++ b/cpp/src/core/coo_sort.cu.in @@ -0,0 +1,65 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include + +#include +#include +#include + +#include + +#include + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, int, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, long, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, int, int*, int*, double*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, long, int*, int*, double*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, unsigned long, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, long, long, unsigned long, long*, long*, float*, cudaStream_t); + +namespace cuvs::detail { + +template +void coo_sort(IdxT m, IdxT n, nnz_t nnz, IdxT* rows, IdxT* cols, T* vals, cudaStream_t stream) + +{ + return raft::sparse::op::coo_sort(m, n, nnz, rows, cols, vals, stream); +} + +template +void coo_sort( + raft::resources const& handle, IdxT m, IdxT n, nnz_t nnz, IdxT* rows, IdxT* cols, T* vals) + +{ + return raft::sparse::op::coo_sort(handle, m, n, nnz, rows, cols, vals); +} + +template +void coo_sort(COO* const in, cudaStream_t stream) + +{ + return raft::sparse::op::coo_sort(in, stream); +} + +} // namespace cuvs::detail + +// clang-format off +@coo_merge_kernel_extern_instantiations@ +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on diff --git a/cpp/src/core/coo_sort_launchers.hpp b/cpp/src/core/coo_sort_launchers.hpp new file mode 100644 index 0000000000..329588508d --- /dev/null +++ b/cpp/src/core/coo_sort_launchers.hpp @@ -0,0 +1,29 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include +#include +#include +#include + +namespace cuvs::detail { +using raft::sparse::COO; +template +void coo_sort(IdxT m, IdxT n, nnz_t nnz, IdxT* rows, IdxT* cols, T* vals, cudaStream_t stream); + +template +void coo_sort( + raft::resources const& handle, IdxT m, IdxT n, nnz_t nnz, IdxT* rows, IdxT* cols, T* vals); + +template +void coo_sort(COO* const in, cudaStream_t stream); + +template +CUVS_EXPORT void compute_duplicates_mask( + value_idx* mask, const value_idx* rows, const value_idx* cols, size_t nnz, cudaStream_t stream); +} // namespace cuvs::detail diff --git a/cpp/src/core/coo_sort_matrix.json b/cpp/src/core/coo_sort_matrix.json new file mode 100644 index 0000000000..e9e8ffb349 --- /dev/null +++ b/cpp/src/core/coo_sort_matrix.json @@ -0,0 +1,55 @@ +[ + { + "return_type": "void", + "qualified_function": "cuvs::detail::coo_sort", + "template_arguments": "", + "function_qualifiers": "", + "_variants": [ + { + "parameter_types": "int, int, int, int*, int*, float*, CUstream_st*", + "instantiation_name": "cuvs_detail_coo_sort_int_int_int_int_ptr_int_ptr_f32_ptr_stream_ptr" + }, + { + "parameter_types": "int, int, long, int*, int*, float*, CUstream_st*", + "instantiation_name": "cuvs_detail_coo_sort_int_int_long_int_ptr_int_ptr_f32_ptr_stream_ptr" + } + ] + }, + { + "return_type": "void", + "qualified_function": "raft::sparse::op::detail::coo_sort", + "function_qualifiers": "", + "_variants": [ + { + "template_arguments": "", + "parameter_types": "bool, int, int, int, int*, int*, float*, cudaStream_t", + "instantiation_name": "raft_sparse_op_detail_coo_sort_f32_int_int" + }, + { + "template_arguments": "", + "parameter_types": "bool, int, int, long, int*, int*, float*, cudaStream_t", + "instantiation_name": "raft_sparse_op_detail_coo_sort_f32_int_long" + }, + { + "template_arguments": "", + "parameter_types": "bool, int, int, int, int*, int*, double*, cudaStream_t", + "instantiation_name": "raft_sparse_op_detail_coo_sort_f64_int_int" + }, + { + "template_arguments": "", + "parameter_types": "bool, int, int, long, int*, int*, double*, cudaStream_t", + "instantiation_name": "raft_sparse_op_detail_coo_sort_f64_int_long" + }, + { + "template_arguments": "", + "parameter_types": "bool, int, int, unsigned long, int*, int*, float*, cudaStream_t", + "instantiation_name": "raft_sparse_op_detail_coo_sort_f32_int_ulong" + }, + { + "template_arguments": "", + "parameter_types": "bool, long, long, unsigned long, long*, long*, float*, cudaStream_t", + "instantiation_name": "raft_sparse_op_detail_coo_sort_f32_long_ulong" + } + ] + } +] diff --git a/cpp/src/core/coo_symmetrize.cu b/cpp/src/core/coo_symmetrize.cu new file mode 100644 index 0000000000..769dc35826 --- /dev/null +++ b/cpp/src/core/coo_symmetrize.cu @@ -0,0 +1,38 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "cub_ptx_query.hpp" + +#include "spectral_launchers.hpp" + +#include +#include +#include + +#include + +namespace cuvs::detail { + +template +void coo_symmetrize(raft::sparse::COO* in, + raft::sparse::COO* out, + Lambda reduction_op, // two-argument reducer + cudaStream_t stream) + +{ + return raft::sparse::linalg::coo_symmetrize(in, out, reduction_op, stream); +} + +template +void coo_symmetrize(raft::resources const& handle, + raft::device_coo_matrix_view in, + raft::device_coo_matrix& out, + Lambda reduction_op) // two-argument reducer + +{ + return raft::sparse::linalg::coo_symmetrize(handle, in, out, reduction_op); +} + +} // namespace cuvs::detail diff --git a/cpp/src/core/faster_dot_on_csr.cu.in b/cpp/src/core/faster_dot_on_csr.cu.in new file mode 100644 index 0000000000..d17ad2032d --- /dev/null +++ b/cpp/src/core/faster_dot_on_csr.cu.in @@ -0,0 +1,11 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on diff --git a/cpp/src/core/faster_dot_on_csr_matrix.json b/cpp/src/core/faster_dot_on_csr_matrix.json new file mode 100644 index 0000000000..e755ff6046 --- /dev/null +++ b/cpp/src/core/faster_dot_on_csr_matrix.json @@ -0,0 +1,18 @@ +[ + { + "return_type": "void", + "qualified_function": "raft::sparse::linalg::detail::faster_dot_on_csr", + "template_arguments": "", + "parameter_types": "raft::resources const&, float*, long, long const*, long const*, __half const*, __half const*, long, long", + "function_qualifiers": "", + "instantiation_name": "f16_i64_f32" + }, + { + "return_type": "void", + "qualified_function": "raft::sparse::linalg::detail::faster_dot_on_csr", + "template_arguments": "", + "parameter_types": "raft::resources const&, float*, long, long const*, long const*, float const*, float const*, long, long", + "function_qualifiers": "", + "instantiation_name": "f32_i64_f32" + } +] diff --git a/cpp/src/core/fill_indices_by_rows.cu.in b/cpp/src/core/fill_indices_by_rows.cu.in new file mode 100644 index 0000000000..4b5de7a4fa --- /dev/null +++ b/cpp/src/core/fill_indices_by_rows.cu.in @@ -0,0 +1,11 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on diff --git a/cpp/src/core/fill_indices_by_rows_matrix.json b/cpp/src/core/fill_indices_by_rows_matrix.json new file mode 100644 index 0000000000..7d385d167f --- /dev/null +++ b/cpp/src/core/fill_indices_by_rows_matrix.json @@ -0,0 +1,10 @@ +[ + { + "return_type": "void", + "qualified_function": "raft::sparse::convert::detail::fill_indices_by_rows", + "template_arguments": "", + "parameter_types": "raft::resources const&, unsigned int const*, long*, long, long, long, long*, long*, long, size_t", + "function_qualifiers": "", + "instantiation_name": "u32_i64_i64_checked" + } +] diff --git a/cpp/src/core/laplacian_normalized.cu.in b/cpp/src/core/laplacian_normalized.cu.in new file mode 100644 index 0000000000..9a9d2185ee --- /dev/null +++ b/cpp/src/core/laplacian_normalized.cu.in @@ -0,0 +1,115 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include + +#include + +#include +#include +#include + +#include + +extern template void raft::linalg::detail:: + map( + cuda::stream_ref, float*, int, raft::sparse::linalg::detail::zero_to_one_functor, float const*); + +extern template void raft::linalg::detail:: + map( + cuda::stream_ref, + double*, + int, + raft::sparse::linalg::detail::zero_to_one_functor, + double const*); + +extern template void raft::linalg::detail::map( + cuda::stream_ref, double*, int, raft::sqrt_op, double const*); + +extern template void raft::linalg::detail::map( + cuda::stream_ref, float*, int, raft::sqrt_op, float const*); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, double*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, int*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, float*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, float*, long, raft::const_op); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, int, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, long, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, int, int*, int*, double*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, long, int*, int*, double*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, unsigned long, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, long, long, unsigned long, long*, long*, float*, cudaStream_t); + +extern template raft::device_coo_matrix +raft::sparse::linalg::detail::compute_graph_laplacian( + raft::resources const&, raft::device_coo_matrix_view); + +extern template raft::device_coo_matrix +raft::sparse::linalg::detail::compute_graph_laplacian( + raft::resources const&, raft::device_coo_matrix_view); + +extern template raft::device_coo_matrix +raft::sparse::linalg::detail::compute_graph_laplacian( + raft::resources const&, raft::device_coo_matrix_view); + +extern template raft::device_coo_matrix +raft::sparse::linalg::detail::compute_graph_laplacian( + raft::resources const&, raft::device_coo_matrix_view); + +namespace cuvs::detail { + +template +raft::device_csr_matrix, + std::remove_const_t, + std::remove_const_t, + std::remove_const_t> +laplacian_normalized( + raft::resources const& res, + raft::device_csr_matrix_view input, + raft::device_vector_view diagonal_out) + +{ + return raft::sparse::linalg::laplacian_normalized(res, input, diagonal_out); +} + +template +raft::device_coo_matrix, + std::remove_const_t, + std::remove_const_t, + std::remove_const_t> +laplacian_normalized(raft::resources const& res, + raft::device_coo_matrix_view input, + raft::device_vector_view diagonal_out) + +{ + return raft::sparse::linalg::laplacian_normalized(res, input, diagonal_out); +} + +} // namespace cuvs::detail + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on diff --git a/cpp/src/core/laplacian_normalized_matrix.json b/cpp/src/core/laplacian_normalized_matrix.json new file mode 100644 index 0000000000..e7beb11543 --- /dev/null +++ b/cpp/src/core/laplacian_normalized_matrix.json @@ -0,0 +1,34 @@ +[ + { + "return_type": "raft::coo_matrix< std::remove_const::type, std::remove_const::type, std::remove_const::type, std::remove_const::type, true, raft::device_container_policy, (raft::SparsityType)0, std::conditional::type, std::remove_const::type, std::remove_const::type, true, raft::device_container_policy>, raft::coordinate_structure_view::type, std::remove_const::type, std::remove_const::type, true>>::type>", + "qualified_function": "cuvs::detail::laplacian_normalized", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::coo_matrix_view, raft::device_vector_view", + "function_qualifiers": "", + "instantiation_name": "coo_matrix_view_f64_int_int_int_true_vector_f64_int_row" + }, + { + "return_type": "raft::coo_matrix< std::remove_const::type, std::remove_const::type, std::remove_const::type, std::remove_const::type, true, raft::device_container_policy, (raft::SparsityType)0, std::conditional::type, std::remove_const::type, std::remove_const::type, true, raft::device_container_policy>, raft::coordinate_structure_view::type, std::remove_const::type, std::remove_const::type, true>>::type>", + "qualified_function": "cuvs::detail::laplacian_normalized", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::coo_matrix_view, raft::device_vector_view", + "function_qualifiers": "", + "instantiation_name": "coo_matrix_view_f64_int_int_long_true_vector_f64_int_row" + }, + { + "return_type": "raft::coo_matrix< std::remove_const::type, std::remove_const::type, std::remove_const::type, std::remove_const::type, true, raft::device_container_policy, (raft::SparsityType)0, std::conditional::type, std::remove_const::type, std::remove_const::type, true, raft::device_container_policy>, raft::coordinate_structure_view::type, std::remove_const::type, std::remove_const::type, true>>::type>", + "qualified_function": "cuvs::detail::laplacian_normalized", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::coo_matrix_view, raft::device_vector_view", + "function_qualifiers": "", + "instantiation_name": "coo_matrix_view_f32_int_int_int_true_vector_f32_int_row" + }, + { + "return_type": "raft::coo_matrix< std::remove_const::type, std::remove_const::type, std::remove_const::type, std::remove_const::type, true, raft::device_container_policy, (raft::SparsityType)0, std::conditional::type, std::remove_const::type, std::remove_const::type, true, raft::device_container_policy>, raft::coordinate_structure_view::type, std::remove_const::type, std::remove_const::type, true>>::type>", + "qualified_function": "cuvs::detail::laplacian_normalized", + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::coo_matrix_view, raft::device_vector_view", + "function_qualifiers": "", + "instantiation_name": "coo_matrix_view_f32_int_int_long_true_vector_f32_int_row" + } +] diff --git a/cpp/src/core/laplacian_reduce_by_key.cu.in b/cpp/src/core/laplacian_reduce_by_key.cu.in new file mode 100644 index 0000000000..8aedb64405 --- /dev/null +++ b/cpp/src/core/laplacian_reduce_by_key.cu.in @@ -0,0 +1,13 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include +#include + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on diff --git a/cpp/src/core/laplacian_reduce_by_key_matrix.json b/cpp/src/core/laplacian_reduce_by_key_matrix.json new file mode 100644 index 0000000000..c41343c0a0 --- /dev/null +++ b/cpp/src/core/laplacian_reduce_by_key_matrix.json @@ -0,0 +1,18 @@ +[ + { + "return_type": "::cuda::std::pair<::cuda::discard_iterator, float*>", + "qualified_function": "thrust::reduce_by_key", + "template_arguments": ", thrust::cuda_cub::execute_on_stream_nosync_base>, int*, float*, ::cuda::discard_iterator, float*>", + "parameter_types": "const thrust::detail::execution_policy_base, thrust::cuda_cub::execute_on_stream_nosync_base>>&, int*, int*, float*, ::cuda::discard_iterator, float*", + "function_qualifiers": "", + "instantiation_name": "f32" + }, + { + "return_type": "::cuda::std::pair<::cuda::discard_iterator, double*>", + "qualified_function": "thrust::reduce_by_key", + "template_arguments": ", thrust::cuda_cub::execute_on_stream_nosync_base>, int*, double*, ::cuda::discard_iterator, double*>", + "parameter_types": "const thrust::detail::execution_policy_base, thrust::cuda_cub::execute_on_stream_nosync_base>>&, int*, int*, double*, ::cuda::discard_iterator, double*", + "function_qualifiers": "", + "instantiation_name": "f64" + } +] diff --git a/cpp/src/core/laplacian_zero_to_one.cu.in b/cpp/src/core/laplacian_zero_to_one.cu.in new file mode 100644 index 0000000000..73ea1aa4f2 --- /dev/null +++ b/cpp/src/core/laplacian_zero_to_one.cu.in @@ -0,0 +1,12 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on diff --git a/cpp/src/core/laplacian_zero_to_one_matrix.json b/cpp/src/core/laplacian_zero_to_one_matrix.json new file mode 100644 index 0000000000..201a536676 --- /dev/null +++ b/cpp/src/core/laplacian_zero_to_one_matrix.json @@ -0,0 +1,18 @@ +[ + { + "return_type": "void", + "qualified_function": "raft::linalg::detail::map", + "template_arguments": "", + "parameter_types": "cuda::stream_ref, float*, int, raft::sparse::linalg::detail::zero_to_one_functor, float const*", + "function_qualifiers": "", + "instantiation_name": "f32_i32" + }, + { + "return_type": "void", + "qualified_function": "raft::linalg::detail::map", + "template_arguments": "", + "parameter_types": "cuda::stream_ref, double*, int, raft::sparse::linalg::detail::zero_to_one_functor, double const*", + "function_qualifiers": "", + "instantiation_name": "f64_i32" + } +] diff --git a/cpp/src/core/masked_matmul.cu.in b/cpp/src/core/masked_matmul.cu.in new file mode 100644 index 0000000000..0b624a08b1 --- /dev/null +++ b/cpp/src/core/masked_matmul.cu.in @@ -0,0 +1,84 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include + +#include + +#include + +#include +#include +#include + +#include + +extern template void raft::sparse::linalg::detail::faster_dot_on_csr( + raft::resources const&, + float*, + long, + long const*, + long const*, + __half const*, + __half const*, + long, + long); + +extern template void raft::sparse::linalg::detail::faster_dot_on_csr( + raft::resources const&, + float*, + long, + long const*, + long const*, + float const*, + float const*, + long, + long); + +extern template void +raft::sparse::convert::detail::fill_indices_by_rows( + raft::resources const&, unsigned int const*, long*, long, long, long, long*, long*, long, size_t); + +extern template void raft::sparse::convert::detail::calc_nnz_by_rows( + raft::resources const&, unsigned int const*, long, long, long*, size_t&, long&); + +extern template void raft::sparse::convert::detail::gpu_repeat_csr( + raft::resources const&, long const*, long const*, long, long, long*, long*); + +namespace cuvs::detail { + +template +void masked_matmul(raft::resources const& handle, + raft::device_matrix_view A, + raft::device_matrix_view B, + raft::core::bitmap_view mask, + raft::device_csr_matrix_view C, + std::optional> alpha, + std::optional> beta) + +{ + return raft::sparse::linalg::masked_matmul(handle, A, B, mask, C, alpha, beta); +} + +template +void masked_matmul(raft::resources const& handle, + raft::device_matrix_view A, + raft::device_matrix_view B, + raft::core::bitset_view mask, + raft::device_csr_matrix_view C, + std::optional> alpha, + std::optional> beta) + +{ + return raft::sparse::linalg::masked_matmul(handle, A, B, mask, C, alpha, beta); +} + +} // namespace cuvs::detail + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on diff --git a/cpp/src/core/masked_matmul_matrix.json b/cpp/src/core/masked_matmul_matrix.json new file mode 100644 index 0000000000..df415d3631 --- /dev/null +++ b/cpp/src/core/masked_matmul_matrix.json @@ -0,0 +1,26 @@ +[ + { + "return_type": "void", + "qualified_function": "cuvs::detail::masked_matmul", + "template_arguments": "", + "function_qualifiers": "", + "_variants": [ + { + "parameter_types": "raft::resources const&, raft::device_matrix_view<__half const, long, raft::row_major>, raft::device_matrix_view<__half const, long, raft::row_major>, raft::core::bitmap_view, raft::csr_matrix_view, std::optional, cuda::std::layout_right, raft::host_device_accessor, (raft::memory_type)0>>>, std::optional, cuda::std::layout_right, raft::host_device_accessor, (raft::memory_type)0>>>", + "instantiation_name": "matrix_f16_c_long_row_bitmap_view_u32_csr_f32_true_optional_mdspan_extents_1ul_layout_right_host_device_accessor_default_memory_type_0" + }, + { + "parameter_types": "raft::resources const&, raft::device_matrix_view<__half const, long, raft::row_major>, raft::device_matrix_view<__half const, long, raft::row_major>, raft::core::bitset_view, raft::csr_matrix_view, std::optional, cuda::std::layout_right, raft::host_device_accessor, (raft::memory_type)0>>>, std::optional, cuda::std::layout_right, raft::host_device_accessor, (raft::memory_type)0>>>", + "instantiation_name": "matrix_f16_c_long_row_bitset_view_u32_csr_f32_true_optional_mdspan_extents_1ul_layout_right_host_device_accessor_default_memory_type_0" + }, + { + "parameter_types": "raft::resources const&, raft::device_matrix_view, raft::device_matrix_view, raft::core::bitmap_view, raft::csr_matrix_view, std::optional, cuda::std::layout_right, raft::host_device_accessor, (raft::memory_type)0>>>, std::optional, cuda::std::layout_right, raft::host_device_accessor, (raft::memory_type)0>>>", + "instantiation_name": "matrix_f32_c_long_row_bitmap_view_u32_csr_true_optional_mdspan_extents_1ul_layout_right_host_device_accessor_default_memory_type_0" + }, + { + "parameter_types": "raft::resources const&, raft::device_matrix_view, raft::device_matrix_view, raft::core::bitset_view, raft::csr_matrix_view, std::optional, cuda::std::layout_right, raft::host_device_accessor, (raft::memory_type)0>>>, std::optional, cuda::std::layout_right, raft::host_device_accessor, (raft::memory_type)0>>>", + "instantiation_name": "matrix_f32_c_long_row_bitset_view_u32_csr_true_optional_mdspan_extents_1ul_layout_right_host_device_accessor_default_memory_type_0" + } + ] + } +] diff --git a/cpp/src/core/spectral_launchers.hpp b/cpp/src/core/spectral_launchers.hpp new file mode 100644 index 0000000000..306f7e87ae --- /dev/null +++ b/cpp/src/core/spectral_launchers.hpp @@ -0,0 +1,129 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace cuvs::detail { +template +raft::device_csr_matrix, + std::remove_const_t, + std::remove_const_t, + std::remove_const_t> +compute_graph_laplacian( + raft::resources const& res, + raft::device_csr_matrix_view input); + +template +raft::device_coo_matrix compute_graph_laplacian( + raft::resources const& res, + raft::device_coo_matrix_view input); + +template +raft::device_csr_matrix, + std::remove_const_t, + std::remove_const_t, + std::remove_const_t> +laplacian_normalized( + raft::resources const& res, + raft::device_csr_matrix_view input, + raft::device_vector_view diagonal_out); + +template +raft::device_coo_matrix, + std::remove_const_t, + std::remove_const_t, + std::remove_const_t> +laplacian_normalized(raft::resources const& res, + raft::device_coo_matrix_view input, + raft::device_vector_view diagonal_out); + +template +CUVS_EXPORT void symmetrize(raft::resources const& handle, + const value_idx* rows, + const value_idx* cols, + const value_t* vals, + value_idx m, + value_idx n, + nnz_t nnz, + raft::sparse::COO& out); + +template +void masked_matmul(raft::resources const& handle, + raft::device_matrix_view A, + raft::device_matrix_view B, + raft::core::bitmap_view mask, + raft::device_csr_matrix_view C, + std::optional> alpha = std::nullopt, + std::optional> beta = std::nullopt); + +template +void masked_matmul(raft::resources const& handle, + raft::device_matrix_view A, + raft::device_matrix_view B, + raft::core::bitset_view mask, + raft::device_csr_matrix_view C, + std::optional> alpha = std::nullopt, + std::optional> beta = std::nullopt); +} // namespace cuvs::detail + +namespace cuvs::detail { +template +void coo_symmetrize(raft::sparse::COO* in, + raft::sparse::COO* out, + Lambda reduction_op, // two-argument reducer + cudaStream_t stream); + +template +void coo_symmetrize(raft::resources const& handle, + raft::device_coo_matrix_view in, + raft::device_coo_matrix& out, + Lambda reduction_op); + +template +void coo_remove_scalar(const idx_t* rows, + const idx_t* cols, + const T* vals, + nnz_t nnz, + idx_t* crows, + idx_t* ccols, + T* cvals, + nnz_t* cnnz, + nnz_t* cur_cnnz, + T scalar, + idx_t n, + cudaStream_t stream); + +template +void coo_remove_scalar(raft::sparse::COO* in, + raft::sparse::COO* out, + T scalar, + cudaStream_t stream); + +template +void coo_remove_scalar(raft::resources const& handle, + raft::device_coo_matrix_view in, + raft::host_scalar_view scalar, + raft::device_coo_matrix& out); +} // namespace cuvs::detail + +namespace cuvs::detail { +template +void symmetrize_connectivity_graph( + raft::resources const& handle, + raft::device_coo_matrix_view input, + raft::device_coo_matrix& output); +} diff --git a/cpp/src/core/symmetrize.cu.in b/cpp/src/core/symmetrize.cu.in new file mode 100644 index 0000000000..f43b56773f --- /dev/null +++ b/cpp/src/core/symmetrize.cu.in @@ -0,0 +1,62 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include + +#include + +#include +#include +#include + +#include + +extern template void raft::sparse::op::detail::compute_duplicates_mask( + int*, int const*, int const*, unsigned long, cudaStream_t); + +extern template void raft::sparse::op::detail::compute_duplicates_mask( + long*, long const*, long const*, unsigned long, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, int, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, long, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, int, int*, int*, double*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, long, int*, int*, double*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, int, int, unsigned long, int*, int*, float*, cudaStream_t); + +extern template void raft::sparse::op::detail::coo_sort( + bool, long, long, unsigned long, long*, long*, float*, cudaStream_t); + +namespace cuvs::detail { + +template +void symmetrize(raft::resources const& handle, + const value_idx* rows, + const value_idx* cols, + const value_t* vals, + value_idx m, + value_idx n, + nnz_t nnz, + raft::sparse::COO& out) + +{ + return raft::sparse::linalg::symmetrize(handle, rows, cols, vals, m, n, nnz, out); +} + +} // namespace cuvs::detail + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on diff --git a/cpp/src/core/symmetrize_connectivity_graph.cu.in b/cpp/src/core/symmetrize_connectivity_graph.cu.in new file mode 100644 index 0000000000..ef87e38347 --- /dev/null +++ b/cpp/src/core/symmetrize_connectivity_graph.cu.in @@ -0,0 +1,85 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include + +#include + +#include + +#include +#include +#include + +#include + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, int*, long, raft::const_op); + +extern template thrust::device_ptr thrust::cuda_cub::exclusive_scan, + thrust::device_ptr, + int, + ::cuda::std::plus>( + thrust::cuda_cub::execution_policy&, + thrust::device_ptr, + thrust::device_ptr, + thrust::device_ptr, + int, + ::cuda::std::plus); + +extern template thrust::device_ptr thrust::cuda_cub::exclusive_scan, + thrust::device_ptr, + long, + ::cuda::std::plus>( + thrust::cuda_cub::execution_policy&, + thrust::device_ptr, + thrust::device_ptr, + thrust::device_ptr, + long, + ::cuda::std::plus); + +extern template void raft::sparse::linalg::detail::coo_degree<64, int, int, int>(int const*, + int, + int*, + cudaStream_t); + +extern template void raft::sparse::linalg::detail::coo_degree<64, int, long, long>(int const*, + long, + long*, + cudaStream_t); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, double*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, int*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, float*, int, raft::const_op); + +extern template void raft::linalg::detail::map>( + cuda::stream_ref, float*, long, raft::const_op); + +namespace cuvs::detail { + +template +void symmetrize_connectivity_graph( + raft::resources const& handle, + raft::device_coo_matrix_view input, + raft::device_coo_matrix& output) +{ + raft::sparse::linalg::coo_symmetrize<128, float, int, NNZType>( + handle, input, output, [] __device__(int row, int col, float a, float b) { + return 0.5f * (a + b); + }); +} + +// clang-format off +template @return_type@ @qualified_function@@template_arguments@(@parameter_types@)@function_qualifiers@; +// clang-format on + +} // namespace cuvs::detail diff --git a/cpp/src/core/symmetrize_connectivity_graph_matrix.json b/cpp/src/core/symmetrize_connectivity_graph_matrix.json new file mode 100644 index 0000000000..9aeda955c0 --- /dev/null +++ b/cpp/src/core/symmetrize_connectivity_graph_matrix.json @@ -0,0 +1,20 @@ +[ + { + "return_type": "void", + "qualified_function": "symmetrize_connectivity_graph", + "function_qualifiers": "", + "namespace_context": "cuvs::detail", + "_variants": [ + { + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::device_coo_matrix_view, raft::device_coo_matrix&", + "instantiation_name": "int" + }, + { + "template_arguments": "", + "parameter_types": "raft::resources const&, raft::device_coo_matrix_view, raft::device_coo_matrix&", + "instantiation_name": "i64" + } + ] + } +] diff --git a/cpp/src/core/symmetrize_matrix.json b/cpp/src/core/symmetrize_matrix.json new file mode 100644 index 0000000000..94255f37e0 --- /dev/null +++ b/cpp/src/core/symmetrize_matrix.json @@ -0,0 +1,18 @@ +[ + { + "return_type": "CUVS_EXPORT void", + "qualified_function": "cuvs::detail::symmetrize", + "template_arguments": "", + "function_qualifiers": "", + "_variants": [ + { + "parameter_types": "raft::resources const&, int const*, int const*, float const*, int, int, unsigned long, raft::sparse::detail::COO&", + "instantiation_name": "int_c_ptr_int_c_ptr_f32_c_ptr_int_int_ulong_coo_f32_int_ulong_ref" + }, + { + "parameter_types": "raft::resources const&, long const*, long const*, float const*, long, long, unsigned long, raft::sparse::detail::COO&", + "instantiation_name": "long_c_ptr_long_c_ptr_f32_c_ptr_long_long_ulong_coo_f32_long_ulong_ref" + } + ] + } +] diff --git a/cpp/src/neighbors/detail/knn_brute_force.cuh b/cpp/src/neighbors/detail/knn_brute_force.cuh index 7952cc09db..ba551ed03a 100644 --- a/cpp/src/neighbors/detail/knn_brute_force.cuh +++ b/cpp/src/neighbors/detail/knn_brute_force.cuh @@ -6,6 +6,7 @@ #pragma once #include "../../core/cub_runtime_launchers.hpp" +#include "../../core/spectral_launchers.hpp" #include #include #include @@ -56,12 +57,6 @@ #include #include -extern template void raft::sparse::convert::detail::calc_nnz_by_rows( - raft::resources const&, unsigned int const*, long, long, long*, size_t&, long&); - -extern template void raft::sparse::convert::detail::gpu_repeat_csr( - raft::resources const&, long const*, long const*, long, long, long*, long*); - namespace cuvs::neighbors::detail { /** @@ -848,7 +843,7 @@ void brute_force_search_filtered( std::visit( [&](const auto& actual_view) { - raft::sparse::linalg::masked_matmul(res, queries, dataset_view, actual_view, csr_view); + cuvs::detail::masked_matmul(res, queries, dataset_view, actual_view, csr_view); }, *filter_view); diff --git a/cpp/src/neighbors/detail/knn_graph.cuh b/cpp/src/neighbors/detail/knn_graph.cuh index ab8e062055..c0b8e5a8ef 100644 --- a/cpp/src/neighbors/detail/knn_graph.cuh +++ b/cpp/src/neighbors/detail/knn_graph.cuh @@ -5,6 +5,7 @@ #pragma once #include "../../core/map_launchers.hpp" +#include "../../core/spectral_launchers.hpp" #include #include @@ -93,13 +94,13 @@ void knn_graph(raft::resources const& res, raft::cast_op{}, raft::make_const_mdspan(indices_64_view)); - raft::sparse::linalg::symmetrize(res, - rows.data(), - indices.data(), - data.data(), - static_cast(m), - static_cast(k), - nnz, - out); + cuvs::detail::symmetrize(res, + rows.data(), + indices.data(), + data.data(), + static_cast(m), + static_cast(k), + nnz, + out); } }; // namespace cuvs::neighbors::detail diff --git a/cpp/src/neighbors/detail/reachability.cuh b/cpp/src/neighbors/detail/reachability.cuh index 338dbd88a2..4afc1828b1 100644 --- a/cpp/src/neighbors/detail/reachability.cuh +++ b/cpp/src/neighbors/detail/reachability.cuh @@ -4,6 +4,7 @@ */ #pragma once +#include "../../core/spectral_launchers.hpp" #include "./knn_brute_force.cuh" #include @@ -215,14 +216,14 @@ void mutual_reachability_graph(const raft::resources& handle, return c / min_samples; }); - raft::sparse::linalg::symmetrize(handle, - coo_rows.data(), - inds.data(), - dists.data(), - m, - m, - static_cast(min_samples * m), - out); + cuvs::detail::symmetrize(handle, + coo_rows.data(), + inds.data(), + dists.data(), + m, + m, + static_cast(min_samples * m), + out); raft::sparse::convert::sorted_coo_to_csr(out.rows(), out.nnz, indptr, m + 1, stream); diff --git a/cpp/src/preprocessing/spectral/detail/spectral_embedding.cuh b/cpp/src/preprocessing/spectral/detail/spectral_embedding.cuh index e1616ef984..1cd4220b67 100644 --- a/cpp/src/preprocessing/spectral/detail/spectral_embedding.cuh +++ b/cpp/src/preprocessing/spectral/detail/spectral_embedding.cuh @@ -5,6 +5,8 @@ #pragma once +#include "../../../core/coo_sort_launchers.hpp" +#include "../../../core/spectral_launchers.hpp" #include #include @@ -34,10 +36,9 @@ OutSparseMatrixType create_laplacian(raft::resources const& handle, InSparseMatrixViewType sparse_matrix_view, raft::device_vector_view diagonal) { - auto laplacian = - spectral_embedding_config.norm_laplacian - ? raft::sparse::linalg::laplacian_normalized(handle, sparse_matrix_view, diagonal) - : raft::sparse::linalg::compute_graph_laplacian(handle, sparse_matrix_view); + auto laplacian = spectral_embedding_config.norm_laplacian + ? cuvs::detail::laplacian_normalized(handle, sparse_matrix_view, diagonal) + : cuvs::detail::compute_graph_laplacian(handle, sparse_matrix_view); auto laplacian_elements_view = raft::make_device_vector_view( laplacian.get_elements().data(), laplacian.structure_view().get_nnz()); @@ -182,21 +183,17 @@ void create_connectivity_graph( auto sym_coo1_matrix = raft::make_device_coo_matrix(handle, n_samples, n_samples); - raft::sparse::linalg::coo_symmetrize<128, float, int, NNZType>( - handle, coo_matrix_view, sym_coo1_matrix, [] __device__(int row, int col, float a, float b) { - return 0.5f * (a + b); - }); - - raft::sparse::op::coo_sort( - n_samples, - n_samples, - sym_coo1_matrix.structure_view().get_nnz(), - sym_coo1_matrix.structure_view().get_rows().data(), - sym_coo1_matrix.structure_view().get_cols().data(), - sym_coo1_matrix.get_elements().data(), - stream.get()); - - raft::sparse::op::coo_remove_scalar<128, float, int, NNZType>( + cuvs::detail::symmetrize_connectivity_graph(handle, coo_matrix_view, sym_coo1_matrix); + + cuvs::detail::coo_sort(n_samples, + n_samples, + sym_coo1_matrix.structure_view().get_nnz(), + sym_coo1_matrix.structure_view().get_rows().data(), + sym_coo1_matrix.structure_view().get_cols().data(), + sym_coo1_matrix.get_elements().data(), + stream.get()); + + cuvs::detail::coo_remove_scalar<128, float, int, NNZType>( handle, raft::make_device_coo_matrix_view( sym_coo1_matrix.get_elements().data(), sym_coo1_matrix.structure_view()), diff --git a/cpp/src/sparse/neighbors/detail/cross_component_nn.cuh b/cpp/src/sparse/neighbors/detail/cross_component_nn.cuh index 65eba7de63..f1bdb266f3 100644 --- a/cpp/src/sparse/neighbors/detail/cross_component_nn.cuh +++ b/cpp/src/sparse/neighbors/detail/cross_component_nn.cuh @@ -8,6 +8,9 @@ #include "../../../core/cub_runtime_launchers.hpp" #include "../../../core/prefix_scans_launchers.hpp" + +#include "../../../core/coo_sort_launchers.hpp" +#include "../../../core/spectral_launchers.hpp" #include "../../../distance/masked_nn.cuh" #include @@ -601,7 +604,7 @@ void cross_component_nn( */ // Compute mask of duplicates rmm::device_uvector out_index(n_rows + 1, stream); - raft::sparse::op::compute_duplicates_mask( + cuvs::detail::compute_duplicates_mask( out_index.data(), colors.data(), nn_colors.data(), n_rows, stream.get()); cuvs::detail::thrust_exclusive_scan(raft::resource::get_thrust_policy(handle), @@ -628,7 +631,7 @@ void cross_component_nn( /** * Symmetrize resulting edge list */ - raft::sparse::linalg::symmetrize( + cuvs::detail::symmetrize( handle, min_edges.rows(), min_edges.cols(), min_edges.vals(), n_rows, n_rows, size, out); }