From 7d766dda2eaee781b8a08b9c1debad73565dffbf Mon Sep 17 00:00:00 2001 From: Mickael Ide Date: Wed, 17 Jan 2024 21:10:37 +0100 Subject: [PATCH] Remove unused sources and headers Signed-off-by: Mickael Ide --- build.sh | 6 +- cpp/CMakeLists.txt | 161 --- cpp/bench/ann/src/raft/raft_cagra_wrapper.h | 10 +- cpp/bench/micro/neighbors/cagra_bench.cuh | 4 +- cpp/include/cuvs/neighbors/ball_cover-ext.cuh | 124 -- cpp/include/cuvs/neighbors/ball_cover-inl.cuh | 395 ------ cpp/include/cuvs/neighbors/ball_cover.cuh | 24 - .../cuvs/neighbors/ball_cover_types.hpp | 169 --- .../cuvs/neighbors/brute_force-ext.cuh | 149 --- .../cuvs/neighbors/brute_force-inl.cuh | 355 ------ cpp/include/cuvs/neighbors/brute_force.cuh | 92 -- .../cuvs/neighbors/brute_force_types.hpp | 283 ----- cpp/include/cuvs/neighbors/cagra.hpp | 109 +- cpp/include/cuvs/neighbors/cagra_types.hpp | 363 ------ .../cuvs/neighbors/epsilon_neighborhood.cuh | 123 -- cpp/include/cuvs/neighbors/ivf_flat-ext.cuh | 206 ---- cpp/include/cuvs/neighbors/ivf_flat-inl.cuh | 602 --------- cpp/include/cuvs/neighbors/ivf_flat.cuh | 24 - .../cuvs/neighbors/ivf_flat_codepacker.hpp | 90 -- .../cuvs/neighbors/ivf_flat_helpers.cuh | 147 --- .../cuvs/neighbors/ivf_flat_serialize.cuh | 154 --- cpp/include/cuvs/neighbors/ivf_flat_types.hpp | 406 ------ cpp/include/cuvs/neighbors/ivf_list.hpp | 196 --- cpp/include/cuvs/neighbors/ivf_list_types.hpp | 79 -- cpp/include/cuvs/neighbors/ivf_pq-ext.cuh | 227 ---- cpp/include/cuvs/neighbors/ivf_pq-inl.cuh | 529 -------- cpp/include/cuvs/neighbors/ivf_pq.cuh | 24 - cpp/include/cuvs/neighbors/ivf_pq_helpers.cuh | 798 ------------ .../cuvs/neighbors/ivf_pq_serialize.cuh | 146 --- cpp/include/cuvs/neighbors/ivf_pq_types.hpp | 580 --------- .../cuvs/neighbors/neighbors_types.hpp | 63 - cpp/include/cuvs/neighbors/nn_descent.cuh | 181 --- .../cuvs/neighbors/nn_descent_types.hpp | 147 --- cpp/include/cuvs/neighbors/refine-ext.cuh | 78 -- cpp/include/cuvs/neighbors/refine-inl.cuh | 104 -- cpp/include/cuvs/neighbors/refine.cuh | 24 - cpp/include/cuvs/neighbors/sample_filter.cuh | 49 - .../cuvs/neighbors/sample_filter_types.hpp | 175 --- .../cuvs/neighbors/specializations.cuh | 22 - cpp/include/cuvs/spatial/knn/ann.cuh | 83 -- cpp/include/cuvs/spatial/knn/ann_common.h | 103 -- cpp/include/cuvs/spatial/knn/ann_types.hpp | 45 - cpp/include/cuvs/spatial/knn/ball_cover.cuh | 70 -- .../cuvs/spatial/knn/ball_cover_types.hpp | 37 - cpp/include/cuvs/spatial/knn/common.hpp | 23 - .../cuvs/spatial/knn/detail/ann_quantized.cuh | 147 --- .../cuvs/spatial/knn/detail/ann_utils.cuh | 576 --------- .../cuvs/spatial/knn/detail/ball_cover.cuh | 549 --------- .../spatial/knn/detail/ball_cover/common.cuh | 73 -- .../knn/detail/ball_cover/registers-ext.cuh | 129 -- .../knn/detail/ball_cover/registers-inl.cuh | 794 ------------ .../knn/detail/ball_cover/registers.cuh | 24 - .../knn/detail/ball_cover/registers_types.cuh | 66 - .../knn/detail/epsilon_neighborhood.cuh | 241 ---- .../spatial/knn/detail/fused_l2_knn-ext.cuh | 74 -- .../spatial/knn/detail/fused_l2_knn-inl.cuh | 1062 ---------------- .../cuvs/spatial/knn/detail/fused_l2_knn.cuh | 24 - .../spatial/knn/detail/haversine_distance.cuh | 143 --- .../cuvs/spatial/knn/detail/processing.cuh | 189 --- .../cuvs/spatial/knn/detail/processing.hpp | 45 - .../cuvs/spatial/knn/epsilon_neighborhood.cuh | 38 - cpp/include/cuvs/spatial/knn/ivf_flat.cuh | 39 - .../cuvs/spatial/knn/ivf_flat_types.hpp | 40 - cpp/include/cuvs/spatial/knn/ivf_pq.cuh | 39 - cpp/include/cuvs/spatial/knn/ivf_pq_types.hpp | 40 - cpp/include/cuvs/spatial/knn/knn.cuh | 231 ---- .../cuvs/spatial/knn/specializations.cuh | 22 - .../cuvs/spatial/knn/specializations/knn.cuh | 22 - cpp/include/cuvs_runtime/cluster/kmeans.hpp | 97 -- .../cuvs_runtime/distance/fused_l2_nn.hpp | 65 - .../distance/pairwise_distance.hpp | 50 - cpp/include/cuvs_runtime/matrix/select_k.hpp | 32 - .../cuvs_runtime/neighbors/brute_force.hpp | 38 - .../cuvs_runtime/neighbors/ivf_flat.hpp | 83 -- cpp/include/cuvs_runtime/neighbors/ivf_pq.hpp | 96 -- cpp/include/cuvs_runtime/neighbors/refine.hpp | 48 - cpp/src/cuvs_runtime/cluster/cluster_cost.cuh | 87 -- .../cluster/cluster_cost_double.cu | 33 - .../cluster/cluster_cost_float.cu | 33 - .../cuvs_runtime/cluster/kmeans_fit_double.cu | 33 - .../cuvs_runtime/cluster/kmeans_fit_float.cu | 33 - .../cluster/kmeans_init_plus_plus_double.cu | 31 - .../cluster/kmeans_init_plus_plus_float.cu | 31 - .../cuvs_runtime/cluster/update_centroids.cuh | 72 -- .../cluster/update_centroids_double.cu | 46 - .../cluster/update_centroids_float.cu | 46 - .../cuvs_runtime/distance/fused_l2_min_arg.cu | 105 -- .../distance/pairwise_distance.cu | 52 - .../matrix/select_k_float_int64_t.cu | 36 - .../brute_force_knn_int64_t_float.cu | 47 - .../cuvs_runtime/neighbors/ivf_flat_build.cu | 62 - .../cuvs_runtime/neighbors/ivf_flat_search.cu | 40 - .../neighbors/ivf_flat_serialize.cu | 65 - cpp/src/cuvs_runtime/neighbors/ivfpq_build.cu | 59 - .../neighbors/ivfpq_deserialize.cu | 31 - .../neighbors/ivfpq_search_float_int64_t.cu | 38 - .../neighbors/ivfpq_search_int8_t_int64_t.cu | 38 - .../neighbors/ivfpq_search_uint8_t_int64_t.cu | 38 - .../cuvs_runtime/neighbors/ivfpq_serialize.cu | 31 - .../neighbors/refine_d_int64_t_float.cu | 33 - .../neighbors/refine_d_int64_t_int8_t.cu | 33 - .../neighbors/refine_d_int64_t_uint8_t.cu | 33 - .../neighbors/refine_h_int64_t_float.cu | 34 - .../neighbors/refine_h_int64_t_int8_t.cu | 33 - .../neighbors/refine_h_int64_t_uint8_t.cu | 33 - cpp/src/cuvs_runtime/random/common.cuh | 41 - ...rmat_rectangular_generator_int64_double.cu | 23 - .../rmat_rectangular_generator_int64_float.cu | 23 - .../rmat_rectangular_generator_int_double.cu | 23 - .../rmat_rectangular_generator_int_float.cu | 23 - .../pairwise_matrix/dispatch_00_generate.py | 194 --- ...patch_canberra_double_double_double_int.cu | 55 - ...dispatch_canberra_float_float_float_int.cu | 50 - ...ch_correlation_double_double_double_int.cu | 55 - ...patch_correlation_float_float_float_int.cu | 55 - ...ispatch_cosine_double_double_double_int.cu | 51 - .../dispatch_cosine_float_float_float_int.cu | 51 - ...ing_unexpanded_double_double_double_int.cu | 50 - ...amming_unexpanded_float_float_float_int.cu | 50 - ...inger_expanded_double_double_double_int.cu | 55 - ...ellinger_expanded_float_float_float_int.cu | 50 - ...jensen_shannon_double_double_double_int.cu | 55 - ...ch_jensen_shannon_float_float_float_int.cu | 55 - ..._kl_divergence_double_double_double_int.cu | 50 - ...tch_kl_divergence_float_float_float_int.cu | 50 - .../dispatch_l1_double_double_double_int.cu | 50 - .../dispatch_l1_float_float_float_int.cu | 50 - ...ch_l2_expanded_double_double_double_int.cu | 51 - ...patch_l2_expanded_float_float_float_int.cu | 51 - ..._l2_unexpanded_double_double_double_int.cu | 55 - ...tch_l2_unexpanded_float_float_float_int.cu | 50 - ...dispatch_l_inf_double_double_double_int.cu | 50 - .../dispatch_l_inf_float_float_float_int.cu | 50 - ..._lp_unexpanded_double_double_double_int.cu | 55 - ...tch_lp_unexpanded_float_float_float_int.cu | 50 - .../detail/pairwise_matrix/dispatch_rbf.cu | 64 - ...tch_russel_rao_double_double_double_int.cu | 55 - ...spatch_russel_rao_float_float_float_int.cu | 50 - cpp/src/distance/distance.cu | 934 -------------- cpp/src/distance/fused_l2_nn.cu | 54 - .../matrix/detail/select_k_double_int64_t.cu | 34 - .../matrix/detail/select_k_double_uint32_t.cu | 35 - cpp/src/matrix/detail/select_k_float_int32.cu | 34 - .../matrix/detail/select_k_float_int64_t.cu | 34 - .../matrix/detail/select_k_float_uint32_t.cu | 34 - .../matrix/detail/select_k_half_int64_t.cu | 34 - .../matrix/detail/select_k_half_uint32_t.cu | 34 - cpp/src/neighbors/ball_cover.cu | 66 - cpp/src/neighbors/brute_force_00_generate.py | 106 -- .../brute_force_fused_l2_knn_float_int64_t.cu | 45 - .../neighbors/brute_force_knn_index_float.cu | 39 - .../brute_force_knn_int64_t_float_int64_t.cu | 47 - .../brute_force_knn_int64_t_float_uint32_t.cu | 47 - .../brute_force_knn_int_float_int.cu | 47 - ...brute_force_knn_uint32_t_float_uint32_t.cu | 47 - cpp/src/neighbors/cagra_build_float.cpp | 68 +- cpp/src/neighbors/cagra_build_int8.cpp | 68 +- cpp/src/neighbors/cagra_build_uint8.cpp | 68 +- cpp/src/neighbors/cagra_optimize.cpp | 22 +- cpp/src/neighbors/cagra_search_float.cpp | 20 +- cpp/src/neighbors/cagra_search_int8.cpp | 20 +- cpp/src/neighbors/cagra_search_uint8.cpp | 20 +- .../cagra/search_multi_cta_00_generate.py | 108 -- ...arch_multi_cta_float_uint32_dim1024_t32.cu | 66 - ...search_multi_cta_float_uint32_dim128_t8.cu | 66 - ...earch_multi_cta_float_uint32_dim256_t16.cu | 66 - ...earch_multi_cta_float_uint32_dim512_t32.cu | 66 - ...arch_multi_cta_float_uint64_dim1024_t32.cu | 66 - ...search_multi_cta_float_uint64_dim128_t8.cu | 66 - ...earch_multi_cta_float_uint64_dim256_t16.cu | 66 - ...earch_multi_cta_float_uint64_dim512_t32.cu | 66 - ...earch_multi_cta_int8_uint32_dim1024_t32.cu | 66 - .../search_multi_cta_int8_uint32_dim128_t8.cu | 66 - ...search_multi_cta_int8_uint32_dim256_t16.cu | 66 - ...search_multi_cta_int8_uint32_dim512_t32.cu | 66 - ...arch_multi_cta_uint8_uint32_dim1024_t32.cu | 66 - ...search_multi_cta_uint8_uint32_dim128_t8.cu | 66 - ...earch_multi_cta_uint8_uint32_dim256_t16.cu | 66 - ...earch_multi_cta_uint8_uint32_dim512_t32.cu | 66 - .../cagra/search_single_cta_00_generate.py | 113 -- ...rch_single_cta_float_uint32_dim1024_t32.cu | 67 - ...earch_single_cta_float_uint32_dim128_t8.cu | 67 - ...arch_single_cta_float_uint32_dim256_t16.cu | 67 - ...arch_single_cta_float_uint32_dim512_t32.cu | 67 - ...rch_single_cta_float_uint64_dim1024_t32.cu | 67 - ...earch_single_cta_float_uint64_dim128_t8.cu | 67 - ...arch_single_cta_float_uint64_dim256_t16.cu | 67 - ...arch_single_cta_float_uint64_dim512_t32.cu | 67 - ...arch_single_cta_int8_uint32_dim1024_t32.cu | 67 - ...search_single_cta_int8_uint32_dim128_t8.cu | 67 - ...earch_single_cta_int8_uint32_dim256_t16.cu | 67 - ...earch_single_cta_int8_uint32_dim512_t32.cu | 67 - ...rch_single_cta_uint8_uint32_dim1024_t32.cu | 67 - ...earch_single_cta_uint8_uint32_dim128_t8.cu | 67 - ...arch_single_cta_uint8_uint32_dim256_t16.cu | 67 - ...arch_single_cta_uint8_uint32_dim512_t32.cu | 67 - ...at_interleaved_scan_float_float_int64_t.cu | 42 - ...interleaved_scan_int8_t_int32_t_int64_t.cu | 42 - ...terleaved_scan_uint8_t_uint32_t_int64_t.cu | 42 - cpp/src/neighbors/detail/ivf_flat_search.cu | 40 - .../ivf_pq_compute_similarity_00_generate.py | 108 -- .../ivf_pq_compute_similarity_float_float.cu | 81 -- ...f_pq_compute_similarity_float_fp8_false.cu | 81 -- ...vf_pq_compute_similarity_float_fp8_true.cu | 81 -- .../ivf_pq_compute_similarity_float_half.cu | 81 -- ...vf_pq_compute_similarity_half_fp8_false.cu | 81 -- ...ivf_pq_compute_similarity_half_fp8_true.cu | 81 -- .../ivf_pq_compute_similarity_half_half.cu | 81 -- .../detail/refine_host_float_float.cpp | 29 - .../detail/refine_host_int8_t_float.cpp | 29 - .../detail/refine_host_uint8_t_float.cpp | 30 - .../detail/selection_faiss_00_generate.py | 79 -- .../detail/selection_faiss_int32_t_float.cu | 44 - .../detail/selection_faiss_int64_t_double.cu | 44 - .../detail/selection_faiss_int64_t_half.cu | 44 - .../detail/selection_faiss_int_double.cu | 44 - .../detail/selection_faiss_long_float.cu | 44 - .../detail/selection_faiss_size_t_double.cu | 44 - .../detail/selection_faiss_size_t_float.cu | 44 - .../detail/selection_faiss_uint32_t_double.cu | 44 - .../detail/selection_faiss_uint32_t_float.cu | 44 - .../detail/selection_faiss_uint32_t_half.cu | 44 - cpp/src/neighbors/ivf_flat_00_generate.py | 148 --- .../neighbors/ivf_flat_build_float_int64_t.cu | 50 - .../ivf_flat_build_int8_t_int64_t.cu | 50 - .../ivf_flat_build_uint8_t_int64_t.cu | 50 - .../ivf_flat_extend_float_int64_t.cu | 58 - .../ivf_flat_extend_int8_t_int64_t.cu | 58 - .../ivf_flat_extend_uint8_t_int64_t.cu | 58 - .../ivf_flat_search_float_int64_t.cu | 49 - .../ivf_flat_search_int8_t_int64_t.cu | 49 - .../ivf_flat_search_uint8_t_int64_t.cu | 49 - .../neighbors/ivfpq_build_float_int64_t.cu | 36 - .../neighbors/ivfpq_build_int8_t_int64_t.cu | 36 - .../neighbors/ivfpq_build_uint8_t_int64_t.cu | 36 - .../neighbors/ivfpq_extend_float_int64_t.cu | 50 - .../neighbors/ivfpq_extend_int8_t_int64_t.cu | 50 - .../neighbors/ivfpq_extend_uint8_t_int64_t.cu | 50 - .../neighbors/ivfpq_search_float_int64_t.cu | 42 - .../neighbors/ivfpq_search_int8_t_int64_t.cu | 42 - .../neighbors/ivfpq_search_uint8_t_int64_t.cu | 42 - cpp/src/neighbors/refine_00_generate.py | 78 -- cpp/src/neighbors/refine_float_float.cu | 50 - cpp/src/neighbors/refine_int8_t_float.cu | 50 - cpp/src/neighbors/refine_uint8_t_float.cu | 50 - .../knn/detail/ball_cover/registers.cu | 60 - .../ball_cover/registers_00_generate.py | 112 -- .../ball_cover/registers_pass_one_2d_dist.cu | 48 - .../registers_pass_one_2d_euclidean.cu | 48 - .../registers_pass_one_2d_haversine.cu | 48 - .../ball_cover/registers_pass_one_3d_dist.cu | 48 - .../registers_pass_one_3d_euclidean.cu | 48 - .../registers_pass_one_3d_haversine.cu | 48 - .../ball_cover/registers_pass_two_2d_dist.cu | 48 - .../registers_pass_two_2d_euclidean.cu | 48 - .../registers_pass_two_2d_haversine.cu | 48 - .../ball_cover/registers_pass_two_3d_dist.cu | 48 - .../registers_pass_two_3d_euclidean.cu | 48 - .../registers_pass_two_3d_haversine.cu | 48 - .../knn/detail/fused_l2_knn_int32_t_float.cu | 42 - .../knn/detail/fused_l2_knn_int64_t_float.cu | 42 - .../knn/detail/fused_l2_knn_uint32_t_float.cu | 43 - cpp/template/src/cagra_example.cu | 2 +- cpp/test/CMakeLists.txt | 124 -- cpp/test/cluster/cluster_solvers.cu | 103 -- cpp/test/cluster/kmeans.cu | 359 ------ cpp/test/cluster/kmeans_balanced.cu | 236 ---- cpp/test/cluster/kmeans_find_k.cu | 140 --- cpp/test/cluster/linkage.cu | 675 ---------- cpp/test/distance/dist_adj.cu | 194 --- cpp/test/distance/dist_adj.cuh | 71 -- .../distance/dist_adj_distance_instance.cu | 63 - cpp/test/distance/dist_adj_threshold.cuh | 36 - cpp/test/distance/dist_canberra.cu | 70 -- cpp/test/distance/dist_correlation.cu | 94 -- cpp/test/distance/dist_cos.cu | 110 -- cpp/test/distance/dist_hamming.cu | 71 -- cpp/test/distance/dist_hellinger.cu | 71 -- cpp/test/distance/dist_inner_product.cu | 74 -- cpp/test/distance/dist_jensen_shannon.cu | 71 -- cpp/test/distance/dist_kl_divergence.cu | 71 -- cpp/test/distance/dist_l1.cu | 70 -- cpp/test/distance/dist_l2_exp.cu | 113 -- cpp/test/distance/dist_l2_sqrt_exp.cu | 74 -- cpp/test/distance/dist_l2_unexp.cu | 71 -- cpp/test/distance/dist_l_inf.cu | 70 -- cpp/test/distance/dist_lp_unexp.cu | 71 -- cpp/test/distance/dist_russell_rao.cu | 71 -- cpp/test/distance/distance_base.cuh | 673 ---------- cpp/test/distance/fused_l2_nn.cu | 436 ------- cpp/test/distance/gram.cu | 170 --- cpp/test/distance/gram_base.cuh | 88 -- cpp/test/distance/masked_nn.cu | 435 ------- .../distance/masked_nn_compress_to_bits.cu | 217 ---- cpp/test/ext_headers/00_generate.py | 79 -- cpp/test/ext_headers/raft_core_logger.cpp | 27 - ...istance_detail_pairwise_matrix_dispatch.cu | 27 - .../ext_headers/raft_distance_distance.cu | 27 - .../ext_headers/raft_distance_fused_l2_nn.cu | 27 - .../raft_linalg_detail_coalesced_reduction.cu | 27 - .../raft_matrix_detail_select_k.cu | 27 - .../ext_headers/raft_neighbors_ball_cover.cu | 27 - .../ext_headers/raft_neighbors_brute_force.cu | 27 - ...ghbors_detail_ivf_flat_interleaved_scan.cu | 27 - .../raft_neighbors_detail_ivf_flat_search.cu | 27 - ...ghbors_detail_ivf_pq_compute_similarity.cu | 27 - .../raft_neighbors_detail_selection_faiss.cu | 27 - .../ext_headers/raft_neighbors_ivf_flat.cu | 27 - cpp/test/ext_headers/raft_neighbors_ivf_pq.cu | 27 - cpp/test/ext_headers/raft_neighbors_refine.cu | 27 - ...spatial_knn_detail_ball_cover_registers.cu | 27 - .../raft_spatial_knn_detail_fused_l2_knn.cu | 27 - .../ext_headers/raft_util_memory_pool.cpp | 27 - cpp/test/neighbors/ann_cagra.cuh | 22 +- .../ann_cagra/test_float_uint32_t.cu | 13 - .../ann_cagra/test_int8_t_uint32_t.cu | 14 +- .../ann_cagra/test_uint8_t_uint32_t.cu | 14 +- cpp/test/neighbors/ann_ivf_flat.cuh | 615 --------- .../ann_ivf_flat/test_filter_float_int64_t.cu | 29 - .../ann_ivf_flat/test_float_int64_t.cu | 32 - .../ann_ivf_flat/test_int8_t_int64_t.cu | 28 - .../ann_ivf_flat/test_uint8_t_int64_t.cu | 28 - cpp/test/neighbors/ann_ivf_pq.cuh | 1095 ----------------- .../ann_ivf_pq/test_filter_float_int64_t.cu | 26 - .../ann_ivf_pq/test_filter_int8_t_int64_t.cu | 27 - .../ann_ivf_pq/test_float_int64_t.cu | 27 - .../ann_ivf_pq/test_float_uint32_t.cu | 37 - .../ann_ivf_pq/test_int8_t_int64_t.cu | 27 - .../ann_ivf_pq/test_uint8_t_int64_t.cu | 27 - cpp/test/neighbors/ann_nn_descent.cuh | 156 --- .../ann_nn_descent/test_float_uint32_t.cu | 28 - .../ann_nn_descent/test_int8_t_uint32_t.cu | 28 - .../ann_nn_descent/test_uint8_t_uint32_t.cu | 28 - cpp/test/neighbors/ball_cover.cu | 372 ------ cpp/test/neighbors/epsilon_neighborhood.cu | 121 -- cpp/test/neighbors/fused_l2_knn.cu | 175 --- cpp/test/neighbors/haversine.cu | 133 -- cpp/test/neighbors/knn.cu | 197 --- cpp/test/neighbors/knn_utils.cuh | 94 -- cpp/test/neighbors/refine.cu | 129 -- cpp/test/neighbors/selection.cu | 499 -------- cpp/test/neighbors/spatial_data.h | 38 - cpp/test/neighbors/tiled_knn.cu | 354 ------ cpp/test/sparse/dist_coo_spmv.cu | 697 ----------- cpp/test/sparse/distance.cu | 853 ------------- cpp/test/sparse/gram.cu | 328 ----- cpp/test/sparse/neighbors/brute_force.cu | 179 --- .../sparse/neighbors/cross_component_nn.cu | 1036 ---------------- cpp/test/sparse/neighbors/knn_graph.cu | 127 -- cpp/test/sparse/spectral_matrix.cu | 83 -- cpp/test/stats/accuracy.cu | 106 -- cpp/test/stats/adjusted_rand_index.cu | 207 ---- cpp/test/stats/completeness_score.cu | 136 -- cpp/test/stats/contingencyMatrix.cu | 167 --- cpp/test/stats/cov.cu | 193 --- cpp/test/stats/dispersion.cu | 132 -- cpp/test/stats/entropy.cu | 123 -- cpp/test/stats/histogram.cu | 318 ----- cpp/test/stats/homogeneity_score.cu | 134 -- cpp/test/stats/information_criterion.cu | 151 --- cpp/test/stats/kl_divergence.cu | 107 -- cpp/test/stats/mean.cu | 149 --- cpp/test/stats/meanvar.cu | 161 --- cpp/test/stats/minmax.cu | 208 ---- cpp/test/stats/mutual_info_score.cu | 162 --- cpp/test/stats/neighborhood_recall.cu | 178 --- cpp/test/stats/r2_score.cu | 114 -- cpp/test/stats/rand_index.cu | 129 -- cpp/test/stats/regression_metrics.cu | 146 --- cpp/test/stats/silhouette_score.cu | 227 ---- cpp/test/stats/stddev.cu | 196 --- cpp/test/stats/sum.cu | 111 -- cpp/test/stats/trustworthiness.cu | 352 ------ cpp/test/stats/v_measure.cu | 139 --- cpp/test/stats/weighted_mean.cu | 339 ----- 375 files changed, 215 insertions(+), 41426 deletions(-) delete mode 100644 cpp/include/cuvs/neighbors/ball_cover-ext.cuh delete mode 100644 cpp/include/cuvs/neighbors/ball_cover-inl.cuh delete mode 100644 cpp/include/cuvs/neighbors/ball_cover.cuh delete mode 100644 cpp/include/cuvs/neighbors/ball_cover_types.hpp delete mode 100644 cpp/include/cuvs/neighbors/brute_force-ext.cuh delete mode 100644 cpp/include/cuvs/neighbors/brute_force-inl.cuh delete mode 100644 cpp/include/cuvs/neighbors/brute_force.cuh delete mode 100644 cpp/include/cuvs/neighbors/brute_force_types.hpp delete mode 100644 cpp/include/cuvs/neighbors/cagra_types.hpp delete mode 100644 cpp/include/cuvs/neighbors/epsilon_neighborhood.cuh delete mode 100644 cpp/include/cuvs/neighbors/ivf_flat-ext.cuh delete mode 100644 cpp/include/cuvs/neighbors/ivf_flat-inl.cuh delete mode 100644 cpp/include/cuvs/neighbors/ivf_flat.cuh delete mode 100644 cpp/include/cuvs/neighbors/ivf_flat_codepacker.hpp delete mode 100644 cpp/include/cuvs/neighbors/ivf_flat_helpers.cuh delete mode 100644 cpp/include/cuvs/neighbors/ivf_flat_serialize.cuh delete mode 100644 cpp/include/cuvs/neighbors/ivf_flat_types.hpp delete mode 100644 cpp/include/cuvs/neighbors/ivf_list.hpp delete mode 100644 cpp/include/cuvs/neighbors/ivf_list_types.hpp delete mode 100644 cpp/include/cuvs/neighbors/ivf_pq-ext.cuh delete mode 100644 cpp/include/cuvs/neighbors/ivf_pq-inl.cuh delete mode 100644 cpp/include/cuvs/neighbors/ivf_pq.cuh delete mode 100644 cpp/include/cuvs/neighbors/ivf_pq_helpers.cuh delete mode 100644 cpp/include/cuvs/neighbors/ivf_pq_serialize.cuh delete mode 100644 cpp/include/cuvs/neighbors/ivf_pq_types.hpp delete mode 100644 cpp/include/cuvs/neighbors/neighbors_types.hpp delete mode 100644 cpp/include/cuvs/neighbors/nn_descent.cuh delete mode 100644 cpp/include/cuvs/neighbors/nn_descent_types.hpp delete mode 100644 cpp/include/cuvs/neighbors/refine-ext.cuh delete mode 100644 cpp/include/cuvs/neighbors/refine-inl.cuh delete mode 100644 cpp/include/cuvs/neighbors/refine.cuh delete mode 100644 cpp/include/cuvs/neighbors/sample_filter.cuh delete mode 100644 cpp/include/cuvs/neighbors/sample_filter_types.hpp delete mode 100644 cpp/include/cuvs/neighbors/specializations.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/ann.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/ann_common.h delete mode 100644 cpp/include/cuvs/spatial/knn/ann_types.hpp delete mode 100644 cpp/include/cuvs/spatial/knn/ball_cover.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/ball_cover_types.hpp delete mode 100644 cpp/include/cuvs/spatial/knn/common.hpp delete mode 100644 cpp/include/cuvs/spatial/knn/detail/ann_quantized.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/ann_utils.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/ball_cover.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/ball_cover/common.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/ball_cover/registers-ext.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/ball_cover/registers-inl.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/ball_cover/registers.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/ball_cover/registers_types.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/epsilon_neighborhood.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/fused_l2_knn-ext.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/fused_l2_knn-inl.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/fused_l2_knn.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/haversine_distance.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/processing.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/detail/processing.hpp delete mode 100644 cpp/include/cuvs/spatial/knn/epsilon_neighborhood.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/ivf_flat.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/ivf_flat_types.hpp delete mode 100644 cpp/include/cuvs/spatial/knn/ivf_pq.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/ivf_pq_types.hpp delete mode 100644 cpp/include/cuvs/spatial/knn/knn.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/specializations.cuh delete mode 100644 cpp/include/cuvs/spatial/knn/specializations/knn.cuh delete mode 100644 cpp/include/cuvs_runtime/cluster/kmeans.hpp delete mode 100644 cpp/include/cuvs_runtime/distance/fused_l2_nn.hpp delete mode 100644 cpp/include/cuvs_runtime/distance/pairwise_distance.hpp delete mode 100644 cpp/include/cuvs_runtime/matrix/select_k.hpp delete mode 100644 cpp/include/cuvs_runtime/neighbors/brute_force.hpp delete mode 100644 cpp/include/cuvs_runtime/neighbors/ivf_flat.hpp delete mode 100644 cpp/include/cuvs_runtime/neighbors/ivf_pq.hpp delete mode 100644 cpp/include/cuvs_runtime/neighbors/refine.hpp delete mode 100644 cpp/src/cuvs_runtime/cluster/cluster_cost.cuh delete mode 100644 cpp/src/cuvs_runtime/cluster/cluster_cost_double.cu delete mode 100644 cpp/src/cuvs_runtime/cluster/cluster_cost_float.cu delete mode 100644 cpp/src/cuvs_runtime/cluster/kmeans_fit_double.cu delete mode 100644 cpp/src/cuvs_runtime/cluster/kmeans_fit_float.cu delete mode 100644 cpp/src/cuvs_runtime/cluster/kmeans_init_plus_plus_double.cu delete mode 100644 cpp/src/cuvs_runtime/cluster/kmeans_init_plus_plus_float.cu delete mode 100644 cpp/src/cuvs_runtime/cluster/update_centroids.cuh delete mode 100644 cpp/src/cuvs_runtime/cluster/update_centroids_double.cu delete mode 100644 cpp/src/cuvs_runtime/cluster/update_centroids_float.cu delete mode 100644 cpp/src/cuvs_runtime/distance/fused_l2_min_arg.cu delete mode 100644 cpp/src/cuvs_runtime/distance/pairwise_distance.cu delete mode 100644 cpp/src/cuvs_runtime/matrix/select_k_float_int64_t.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/brute_force_knn_int64_t_float.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/ivf_flat_build.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/ivf_flat_search.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/ivf_flat_serialize.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/ivfpq_build.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/ivfpq_deserialize.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/ivfpq_search_float_int64_t.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/ivfpq_search_int8_t_int64_t.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/ivfpq_search_uint8_t_int64_t.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/ivfpq_serialize.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/refine_d_int64_t_float.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/refine_d_int64_t_int8_t.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/refine_d_int64_t_uint8_t.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/refine_h_int64_t_float.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/refine_h_int64_t_int8_t.cu delete mode 100644 cpp/src/cuvs_runtime/neighbors/refine_h_int64_t_uint8_t.cu delete mode 100644 cpp/src/cuvs_runtime/random/common.cuh delete mode 100644 cpp/src/cuvs_runtime/random/rmat_rectangular_generator_int64_double.cu delete mode 100644 cpp/src/cuvs_runtime/random/rmat_rectangular_generator_int64_float.cu delete mode 100644 cpp/src/cuvs_runtime/random/rmat_rectangular_generator_int_double.cu delete mode 100644 cpp/src/cuvs_runtime/random/rmat_rectangular_generator_int_float.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_00_generate.py delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_canberra_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_canberra_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_correlation_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_correlation_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_cosine_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_cosine_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_hamming_unexpanded_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_hamming_unexpanded_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_hellinger_expanded_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_hellinger_expanded_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_jensen_shannon_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_jensen_shannon_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_kl_divergence_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_kl_divergence_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_l1_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_l1_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_l2_expanded_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_l2_expanded_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_l2_unexpanded_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_l2_unexpanded_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_l_inf_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_l_inf_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_lp_unexpanded_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_lp_unexpanded_float_float_float_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_rbf.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_russel_rao_double_double_double_int.cu delete mode 100644 cpp/src/distance/detail/pairwise_matrix/dispatch_russel_rao_float_float_float_int.cu delete mode 100644 cpp/src/distance/distance.cu delete mode 100644 cpp/src/distance/fused_l2_nn.cu delete mode 100644 cpp/src/matrix/detail/select_k_double_int64_t.cu delete mode 100644 cpp/src/matrix/detail/select_k_double_uint32_t.cu delete mode 100644 cpp/src/matrix/detail/select_k_float_int32.cu delete mode 100644 cpp/src/matrix/detail/select_k_float_int64_t.cu delete mode 100644 cpp/src/matrix/detail/select_k_float_uint32_t.cu delete mode 100644 cpp/src/matrix/detail/select_k_half_int64_t.cu delete mode 100644 cpp/src/matrix/detail/select_k_half_uint32_t.cu delete mode 100644 cpp/src/neighbors/ball_cover.cu delete mode 100644 cpp/src/neighbors/brute_force_00_generate.py delete mode 100644 cpp/src/neighbors/brute_force_fused_l2_knn_float_int64_t.cu delete mode 100644 cpp/src/neighbors/brute_force_knn_index_float.cu delete mode 100644 cpp/src/neighbors/brute_force_knn_int64_t_float_int64_t.cu delete mode 100644 cpp/src/neighbors/brute_force_knn_int64_t_float_uint32_t.cu delete mode 100644 cpp/src/neighbors/brute_force_knn_int_float_int.cu delete mode 100644 cpp/src/neighbors/brute_force_knn_uint32_t_float_uint32_t.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_00_generate.py delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim1024_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim128_t8.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim256_t16.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim512_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_00_generate.py delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim128_t8.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim256_t16.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim512_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu delete mode 100644 cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu delete mode 100644 cpp/src/neighbors/detail/ivf_flat_interleaved_scan_float_float_int64_t.cu delete mode 100644 cpp/src/neighbors/detail/ivf_flat_interleaved_scan_int8_t_int32_t_int64_t.cu delete mode 100644 cpp/src/neighbors/detail/ivf_flat_interleaved_scan_uint8_t_uint32_t_int64_t.cu delete mode 100644 cpp/src/neighbors/detail/ivf_flat_search.cu delete mode 100644 cpp/src/neighbors/detail/ivf_pq_compute_similarity_00_generate.py delete mode 100644 cpp/src/neighbors/detail/ivf_pq_compute_similarity_float_float.cu delete mode 100644 cpp/src/neighbors/detail/ivf_pq_compute_similarity_float_fp8_false.cu delete mode 100644 cpp/src/neighbors/detail/ivf_pq_compute_similarity_float_fp8_true.cu delete mode 100644 cpp/src/neighbors/detail/ivf_pq_compute_similarity_float_half.cu delete mode 100644 cpp/src/neighbors/detail/ivf_pq_compute_similarity_half_fp8_false.cu delete mode 100644 cpp/src/neighbors/detail/ivf_pq_compute_similarity_half_fp8_true.cu delete mode 100644 cpp/src/neighbors/detail/ivf_pq_compute_similarity_half_half.cu delete mode 100644 cpp/src/neighbors/detail/refine_host_float_float.cpp delete mode 100644 cpp/src/neighbors/detail/refine_host_int8_t_float.cpp delete mode 100644 cpp/src/neighbors/detail/refine_host_uint8_t_float.cpp delete mode 100644 cpp/src/neighbors/detail/selection_faiss_00_generate.py delete mode 100644 cpp/src/neighbors/detail/selection_faiss_int32_t_float.cu delete mode 100644 cpp/src/neighbors/detail/selection_faiss_int64_t_double.cu delete mode 100644 cpp/src/neighbors/detail/selection_faiss_int64_t_half.cu delete mode 100644 cpp/src/neighbors/detail/selection_faiss_int_double.cu delete mode 100644 cpp/src/neighbors/detail/selection_faiss_long_float.cu delete mode 100644 cpp/src/neighbors/detail/selection_faiss_size_t_double.cu delete mode 100644 cpp/src/neighbors/detail/selection_faiss_size_t_float.cu delete mode 100644 cpp/src/neighbors/detail/selection_faiss_uint32_t_double.cu delete mode 100644 cpp/src/neighbors/detail/selection_faiss_uint32_t_float.cu delete mode 100644 cpp/src/neighbors/detail/selection_faiss_uint32_t_half.cu delete mode 100644 cpp/src/neighbors/ivf_flat_00_generate.py delete mode 100644 cpp/src/neighbors/ivf_flat_build_float_int64_t.cu delete mode 100644 cpp/src/neighbors/ivf_flat_build_int8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/ivf_flat_build_uint8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/ivf_flat_extend_float_int64_t.cu delete mode 100644 cpp/src/neighbors/ivf_flat_extend_int8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/ivf_flat_extend_uint8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/ivf_flat_search_float_int64_t.cu delete mode 100644 cpp/src/neighbors/ivf_flat_search_int8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/ivf_flat_search_uint8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/ivfpq_build_float_int64_t.cu delete mode 100644 cpp/src/neighbors/ivfpq_build_int8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/ivfpq_build_uint8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/ivfpq_extend_float_int64_t.cu delete mode 100644 cpp/src/neighbors/ivfpq_extend_int8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/ivfpq_extend_uint8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/ivfpq_search_float_int64_t.cu delete mode 100644 cpp/src/neighbors/ivfpq_search_int8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/ivfpq_search_uint8_t_int64_t.cu delete mode 100644 cpp/src/neighbors/refine_00_generate.py delete mode 100644 cpp/src/neighbors/refine_float_float.cu delete mode 100644 cpp/src/neighbors/refine_int8_t_float.cu delete mode 100644 cpp/src/neighbors/refine_uint8_t_float.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_00_generate.py delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_one_2d_dist.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_one_2d_euclidean.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_one_2d_haversine.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_one_3d_dist.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_one_3d_euclidean.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_one_3d_haversine.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_two_2d_dist.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_two_2d_euclidean.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_two_2d_haversine.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_two_3d_dist.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_two_3d_euclidean.cu delete mode 100644 cpp/src/spatial/knn/detail/ball_cover/registers_pass_two_3d_haversine.cu delete mode 100644 cpp/src/spatial/knn/detail/fused_l2_knn_int32_t_float.cu delete mode 100644 cpp/src/spatial/knn/detail/fused_l2_knn_int64_t_float.cu delete mode 100644 cpp/src/spatial/knn/detail/fused_l2_knn_uint32_t_float.cu delete mode 100644 cpp/test/cluster/cluster_solvers.cu delete mode 100644 cpp/test/cluster/kmeans.cu delete mode 100644 cpp/test/cluster/kmeans_balanced.cu delete mode 100644 cpp/test/cluster/kmeans_find_k.cu delete mode 100644 cpp/test/cluster/linkage.cu delete mode 100644 cpp/test/distance/dist_adj.cu delete mode 100644 cpp/test/distance/dist_adj.cuh delete mode 100644 cpp/test/distance/dist_adj_distance_instance.cu delete mode 100644 cpp/test/distance/dist_adj_threshold.cuh delete mode 100644 cpp/test/distance/dist_canberra.cu delete mode 100644 cpp/test/distance/dist_correlation.cu delete mode 100644 cpp/test/distance/dist_cos.cu delete mode 100644 cpp/test/distance/dist_hamming.cu delete mode 100644 cpp/test/distance/dist_hellinger.cu delete mode 100644 cpp/test/distance/dist_inner_product.cu delete mode 100644 cpp/test/distance/dist_jensen_shannon.cu delete mode 100644 cpp/test/distance/dist_kl_divergence.cu delete mode 100644 cpp/test/distance/dist_l1.cu delete mode 100644 cpp/test/distance/dist_l2_exp.cu delete mode 100644 cpp/test/distance/dist_l2_sqrt_exp.cu delete mode 100644 cpp/test/distance/dist_l2_unexp.cu delete mode 100644 cpp/test/distance/dist_l_inf.cu delete mode 100644 cpp/test/distance/dist_lp_unexp.cu delete mode 100644 cpp/test/distance/dist_russell_rao.cu delete mode 100644 cpp/test/distance/distance_base.cuh delete mode 100644 cpp/test/distance/fused_l2_nn.cu delete mode 100644 cpp/test/distance/gram.cu delete mode 100644 cpp/test/distance/gram_base.cuh delete mode 100644 cpp/test/distance/masked_nn.cu delete mode 100644 cpp/test/distance/masked_nn_compress_to_bits.cu delete mode 100644 cpp/test/ext_headers/00_generate.py delete mode 100644 cpp/test/ext_headers/raft_core_logger.cpp delete mode 100644 cpp/test/ext_headers/raft_distance_detail_pairwise_matrix_dispatch.cu delete mode 100644 cpp/test/ext_headers/raft_distance_distance.cu delete mode 100644 cpp/test/ext_headers/raft_distance_fused_l2_nn.cu delete mode 100644 cpp/test/ext_headers/raft_linalg_detail_coalesced_reduction.cu delete mode 100644 cpp/test/ext_headers/raft_matrix_detail_select_k.cu delete mode 100644 cpp/test/ext_headers/raft_neighbors_ball_cover.cu delete mode 100644 cpp/test/ext_headers/raft_neighbors_brute_force.cu delete mode 100644 cpp/test/ext_headers/raft_neighbors_detail_ivf_flat_interleaved_scan.cu delete mode 100644 cpp/test/ext_headers/raft_neighbors_detail_ivf_flat_search.cu delete mode 100644 cpp/test/ext_headers/raft_neighbors_detail_ivf_pq_compute_similarity.cu delete mode 100644 cpp/test/ext_headers/raft_neighbors_detail_selection_faiss.cu delete mode 100644 cpp/test/ext_headers/raft_neighbors_ivf_flat.cu delete mode 100644 cpp/test/ext_headers/raft_neighbors_ivf_pq.cu delete mode 100644 cpp/test/ext_headers/raft_neighbors_refine.cu delete mode 100644 cpp/test/ext_headers/raft_spatial_knn_detail_ball_cover_registers.cu delete mode 100644 cpp/test/ext_headers/raft_spatial_knn_detail_fused_l2_knn.cu delete mode 100644 cpp/test/ext_headers/raft_util_memory_pool.cpp delete mode 100644 cpp/test/neighbors/ann_ivf_flat.cuh delete mode 100644 cpp/test/neighbors/ann_ivf_flat/test_filter_float_int64_t.cu delete mode 100644 cpp/test/neighbors/ann_ivf_flat/test_float_int64_t.cu delete mode 100644 cpp/test/neighbors/ann_ivf_flat/test_int8_t_int64_t.cu delete mode 100644 cpp/test/neighbors/ann_ivf_flat/test_uint8_t_int64_t.cu delete mode 100644 cpp/test/neighbors/ann_ivf_pq.cuh delete mode 100644 cpp/test/neighbors/ann_ivf_pq/test_filter_float_int64_t.cu delete mode 100644 cpp/test/neighbors/ann_ivf_pq/test_filter_int8_t_int64_t.cu delete mode 100644 cpp/test/neighbors/ann_ivf_pq/test_float_int64_t.cu delete mode 100644 cpp/test/neighbors/ann_ivf_pq/test_float_uint32_t.cu delete mode 100644 cpp/test/neighbors/ann_ivf_pq/test_int8_t_int64_t.cu delete mode 100644 cpp/test/neighbors/ann_ivf_pq/test_uint8_t_int64_t.cu delete mode 100644 cpp/test/neighbors/ann_nn_descent.cuh delete mode 100644 cpp/test/neighbors/ann_nn_descent/test_float_uint32_t.cu delete mode 100644 cpp/test/neighbors/ann_nn_descent/test_int8_t_uint32_t.cu delete mode 100644 cpp/test/neighbors/ann_nn_descent/test_uint8_t_uint32_t.cu delete mode 100644 cpp/test/neighbors/ball_cover.cu delete mode 100644 cpp/test/neighbors/epsilon_neighborhood.cu delete mode 100644 cpp/test/neighbors/fused_l2_knn.cu delete mode 100644 cpp/test/neighbors/haversine.cu delete mode 100644 cpp/test/neighbors/knn.cu delete mode 100644 cpp/test/neighbors/knn_utils.cuh delete mode 100644 cpp/test/neighbors/refine.cu delete mode 100644 cpp/test/neighbors/selection.cu delete mode 100644 cpp/test/neighbors/spatial_data.h delete mode 100644 cpp/test/neighbors/tiled_knn.cu delete mode 100644 cpp/test/sparse/dist_coo_spmv.cu delete mode 100644 cpp/test/sparse/distance.cu delete mode 100644 cpp/test/sparse/gram.cu delete mode 100644 cpp/test/sparse/neighbors/brute_force.cu delete mode 100644 cpp/test/sparse/neighbors/cross_component_nn.cu delete mode 100644 cpp/test/sparse/neighbors/knn_graph.cu delete mode 100644 cpp/test/sparse/spectral_matrix.cu delete mode 100644 cpp/test/stats/accuracy.cu delete mode 100644 cpp/test/stats/adjusted_rand_index.cu delete mode 100644 cpp/test/stats/completeness_score.cu delete mode 100644 cpp/test/stats/contingencyMatrix.cu delete mode 100644 cpp/test/stats/cov.cu delete mode 100644 cpp/test/stats/dispersion.cu delete mode 100644 cpp/test/stats/entropy.cu delete mode 100644 cpp/test/stats/histogram.cu delete mode 100644 cpp/test/stats/homogeneity_score.cu delete mode 100644 cpp/test/stats/information_criterion.cu delete mode 100644 cpp/test/stats/kl_divergence.cu delete mode 100644 cpp/test/stats/mean.cu delete mode 100644 cpp/test/stats/meanvar.cu delete mode 100644 cpp/test/stats/minmax.cu delete mode 100644 cpp/test/stats/mutual_info_score.cu delete mode 100644 cpp/test/stats/neighborhood_recall.cu delete mode 100644 cpp/test/stats/r2_score.cu delete mode 100644 cpp/test/stats/rand_index.cu delete mode 100644 cpp/test/stats/regression_metrics.cu delete mode 100644 cpp/test/stats/silhouette_score.cu delete mode 100644 cpp/test/stats/stddev.cu delete mode 100644 cpp/test/stats/sum.cu delete mode 100644 cpp/test/stats/trustworthiness.cu delete mode 100644 cpp/test/stats/v_measure.cu delete mode 100644 cpp/test/stats/weighted_mean.cu diff --git a/build.sh b/build.sh index eb360ff32..c4b7a7bf7 100755 --- a/build.sh +++ b/build.sh @@ -1,6 +1,6 @@ #!/bin/bash -# Copyright (c) 2020-2023, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. # cuvs build scripts @@ -77,8 +77,8 @@ INSTALL_TARGET=install BUILD_REPORT_METRICS="" BUILD_REPORT_INCL_CACHE_STATS=OFF -TEST_TARGETS="CLUSTER_TEST;DISTANCE_TEST;NEIGHBORS_TEST;NEIGHBORS_ANN_CAGRA_TEST;NEIGHBORS_ANN_NN_DESCENT_TEST;NEIGHBORS_ANN_IVF_TEST" -BENCH_TARGETS="CLUSTER_BENCH;NEIGHBORS_BENCH;DISTANCE_BENCH" +TEST_TARGETS="NEIGHBORS_ANN_CAGRA_TEST" +BENCH_TARGETS="NEIGHBORS_BENCH" CACHE_ARGS="" NVTX=ON diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 8f914227a..2239a7e15 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -191,167 +191,6 @@ include(cmake/thirdparty/get_cutlass.cmake) add_library( cuvs SHARED - # src/distance/detail/pairwise_matrix/dispatch_canberra_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_canberra_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_correlation_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_correlation_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_cosine_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_cosine_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_hamming_unexpanded_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_hamming_unexpanded_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_hellinger_expanded_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_hellinger_expanded_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_jensen_shannon_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_jensen_shannon_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_kl_divergence_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_kl_divergence_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_l1_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_l1_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_l2_expanded_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_l2_expanded_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_l2_unexpanded_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_l2_unexpanded_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_l_inf_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_l_inf_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_lp_unexpanded_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_lp_unexpanded_float_float_float_int.cu - # src/distance/detail/pairwise_matrix/dispatch_rbf.cu - # src/distance/detail/pairwise_matrix/dispatch_russel_rao_double_double_double_int.cu - # src/distance/detail/pairwise_matrix/dispatch_russel_rao_float_float_float_int.cu - # src/distance/distance.cu - # src/distance/fused_l2_nn.cu - # src/matrix/detail/select_k_double_int64_t.cu - # src/matrix/detail/select_k_double_uint32_t.cu - # src/matrix/detail/select_k_float_int64_t.cu - # src/matrix/detail/select_k_float_uint32_t.cu - # src/matrix/detail/select_k_float_int32.cu - # src/matrix/detail/select_k_half_int64_t.cu - # src/matrix/detail/select_k_half_uint32_t.cu - # src/neighbors/ball_cover.cu - # src/neighbors/brute_force_fused_l2_knn_float_int64_t.cu - # src/neighbors/brute_force_knn_int64_t_float_int64_t.cu - # src/neighbors/brute_force_knn_int64_t_float_uint32_t.cu - # src/neighbors/brute_force_knn_int_float_int.cu - # src/neighbors/brute_force_knn_uint32_t_float_uint32_t.cu - # src/neighbors/brute_force_knn_index_float.cu - # src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu - # src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu - # src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu - # src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu - # src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu - # src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu - # src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu - # src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu - # src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu - # src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu - # src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu - # src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu - # src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu - # src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu - # src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu - # src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu - # src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu - # src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu - # src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu - # src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu - # src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu - # src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu - # src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu - # src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu - # src/neighbors/detail/ivf_flat_interleaved_scan_float_float_int64_t.cu - # src/neighbors/detail/ivf_flat_interleaved_scan_int8_t_int32_t_int64_t.cu - # src/neighbors/detail/ivf_flat_interleaved_scan_uint8_t_uint32_t_int64_t.cu - # src/neighbors/detail/ivf_flat_search.cu - # src/neighbors/detail/ivf_pq_compute_similarity_float_float.cu - # src/neighbors/detail/ivf_pq_compute_similarity_float_fp8_false.cu - # src/neighbors/detail/ivf_pq_compute_similarity_float_fp8_true.cu - # src/neighbors/detail/ivf_pq_compute_similarity_float_half.cu - # src/neighbors/detail/ivf_pq_compute_similarity_half_fp8_false.cu - # src/neighbors/detail/ivf_pq_compute_similarity_half_fp8_true.cu - # src/neighbors/detail/ivf_pq_compute_similarity_half_half.cu - # src/neighbors/detail/refine_host_float_float.cpp - # src/neighbors/detail/refine_host_int8_t_float.cpp - # src/neighbors/detail/refine_host_uint8_t_float.cpp - # src/neighbors/detail/selection_faiss_int32_t_float.cu - # src/neighbors/detail/selection_faiss_int_double.cu - # src/neighbors/detail/selection_faiss_long_float.cu - # src/neighbors/detail/selection_faiss_size_t_double.cu - # src/neighbors/detail/selection_faiss_size_t_float.cu - # src/neighbors/detail/selection_faiss_uint32_t_float.cu - # src/neighbors/detail/selection_faiss_int64_t_double.cu - # src/neighbors/detail/selection_faiss_int64_t_half.cu - # src/neighbors/detail/selection_faiss_uint32_t_double.cu - # src/neighbors/detail/selection_faiss_uint32_t_half.cu - # src/neighbors/ivf_flat_build_float_int64_t.cu - # src/neighbors/ivf_flat_build_int8_t_int64_t.cu - # src/neighbors/ivf_flat_build_uint8_t_int64_t.cu - # src/neighbors/ivf_flat_extend_float_int64_t.cu - # src/neighbors/ivf_flat_extend_int8_t_int64_t.cu - # src/neighbors/ivf_flat_extend_uint8_t_int64_t.cu - # src/neighbors/ivf_flat_search_float_int64_t.cu - # src/neighbors/ivf_flat_search_int8_t_int64_t.cu - # src/neighbors/ivf_flat_search_uint8_t_int64_t.cu - # src/neighbors/ivfpq_build_float_int64_t.cu - # src/neighbors/ivfpq_build_int8_t_int64_t.cu - # src/neighbors/ivfpq_build_uint8_t_int64_t.cu - # src/neighbors/ivfpq_extend_float_int64_t.cu - # src/neighbors/ivfpq_extend_int8_t_int64_t.cu - # src/neighbors/ivfpq_extend_uint8_t_int64_t.cu - # src/neighbors/ivfpq_search_float_int64_t.cu - # src/neighbors/ivfpq_search_int8_t_int64_t.cu - # src/neighbors/ivfpq_search_uint8_t_int64_t.cu - # src/neighbors/refine_float_float.cu - # src/neighbors/refine_int8_t_float.cu - # src/neighbors/refine_uint8_t_float.cu - # src/cuvs_runtime/cluster/cluster_cost.cuh - # src/cuvs_runtime/cluster/cluster_cost_double.cu - # src/cuvs_runtime/cluster/cluster_cost_float.cu - # src/cuvs_runtime/cluster/kmeans_fit_double.cu - # src/cuvs_runtime/cluster/kmeans_fit_float.cu - # src/cuvs_runtime/cluster/kmeans_init_plus_plus_double.cu - # src/cuvs_runtime/cluster/kmeans_init_plus_plus_float.cu - # src/cuvs_runtime/cluster/update_centroids.cuh - # src/cuvs_runtime/cluster/update_centroids_double.cu - # src/cuvs_runtime/cluster/update_centroids_float.cu - # src/cuvs_runtime/distance/fused_l2_min_arg.cu - # src/cuvs_runtime/distance/pairwise_distance.cu - # src/cuvs_runtime/matrix/select_k_float_int64_t.cu - # src/cuvs_runtime/neighbors/brute_force_knn_int64_t_float.cu - # src/cuvs_runtime/neighbors/ivf_flat_build.cu - # src/cuvs_runtime/neighbors/ivf_flat_search.cu - # src/cuvs_runtime/neighbors/ivf_flat_serialize.cu - # src/cuvs_runtime/neighbors/ivfpq_build.cu - # src/cuvs_runtime/neighbors/ivfpq_deserialize.cu - # src/cuvs_runtime/neighbors/ivfpq_search_float_int64_t.cu - # src/cuvs_runtime/neighbors/ivfpq_search_int8_t_int64_t.cu - # src/cuvs_runtime/neighbors/ivfpq_search_uint8_t_int64_t.cu - # src/cuvs_runtime/neighbors/ivfpq_serialize.cu - # src/cuvs_runtime/neighbors/refine_d_int64_t_float.cu - # src/cuvs_runtime/neighbors/refine_d_int64_t_int8_t.cu - # src/cuvs_runtime/neighbors/refine_d_int64_t_uint8_t.cu - # src/cuvs_runtime/neighbors/refine_h_int64_t_float.cu - # src/cuvs_runtime/neighbors/refine_h_int64_t_int8_t.cu - # src/cuvs_runtime/neighbors/refine_h_int64_t_uint8_t.cu - # src/cuvs_runtime/random/rmat_rectangular_generator_int64_double.cu - # src/cuvs_runtime/random/rmat_rectangular_generator_int64_float.cu - # src/cuvs_runtime/random/rmat_rectangular_generator_int_double.cu - # src/cuvs_runtime/random/rmat_rectangular_generator_int_float.cu - # src/spatial/knn/detail/ball_cover/registers_pass_one_2d_dist.cu - # src/spatial/knn/detail/ball_cover/registers_pass_one_2d_euclidean.cu - # src/spatial/knn/detail/ball_cover/registers_pass_one_2d_haversine.cu - # src/spatial/knn/detail/ball_cover/registers_pass_one_3d_dist.cu - # src/spatial/knn/detail/ball_cover/registers_pass_one_3d_euclidean.cu - # src/spatial/knn/detail/ball_cover/registers_pass_one_3d_haversine.cu - # src/spatial/knn/detail/ball_cover/registers_pass_two_2d_dist.cu - # src/spatial/knn/detail/ball_cover/registers_pass_two_2d_euclidean.cu - # src/spatial/knn/detail/ball_cover/registers_pass_two_2d_haversine.cu - # src/spatial/knn/detail/ball_cover/registers_pass_two_3d_dist.cu - # src/spatial/knn/detail/ball_cover/registers_pass_two_3d_euclidean.cu - # src/spatial/knn/detail/ball_cover/registers_pass_two_3d_haversine.cu - # src/spatial/knn/detail/fused_l2_knn_int32_t_float.cu - # src/spatial/knn/detail/fused_l2_knn_int64_t_float.cu - # src/spatial/knn/detail/fused_l2_knn_uint32_t_float.cu - src/neighbors/cagra_build_float.cpp src/neighbors/cagra_build_int8.cpp src/neighbors/cagra_build_uint8.cpp diff --git a/cpp/bench/ann/src/raft/raft_cagra_wrapper.h b/cpp/bench/ann/src/raft/raft_cagra_wrapper.h index 90c4218c5..c846416a4 100644 --- a/cpp/bench/ann/src/raft/raft_cagra_wrapper.h +++ b/cpp/bench/ann/src/raft/raft_cagra_wrapper.h @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * Copyright (c) 2023-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -17,9 +17,7 @@ #include #include -#include -#include -#include +#include #include #include #include @@ -57,7 +55,7 @@ class RaftCagra : public ANN { using typename ANN::AnnSearchParam; struct SearchParam : public AnnSearchParam { - cuvs::neighbors::experimental::cagra::search_params p; + cuvs::neighbors::cagra::search_params p; AllocatorType graph_mem = AllocatorType::Device; AllocatorType dataset_mem = AllocatorType::Device; auto needs_dataset() const -> bool override { return true; } @@ -209,7 +207,7 @@ void RaftCagra::set_search_param(const AnnSearchParam& param) allocator_to_string(dataset_mem_).c_str()); auto mr = get_mr(dataset_mem_); - cuvs::neighbors::cagra::detail::copy_with_padding(handle_, dataset_, input_dataset_v_, mr); + raft::neighbors::cagra::detail::copy_with_padding(handle_, dataset_, input_dataset_v_, mr); index_->update_dataset(handle_, make_const_mdspan(dataset_.view())); diff --git a/cpp/bench/micro/neighbors/cagra_bench.cuh b/cpp/bench/micro/neighbors/cagra_bench.cuh index 3be664db8..0cc8c9578 100644 --- a/cpp/bench/micro/neighbors/cagra_bench.cuh +++ b/cpp/bench/micro/neighbors/cagra_bench.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * Copyright (c) 2023-2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -17,7 +17,7 @@ #pragma once #include -#include +#include #include #include #include diff --git a/cpp/include/cuvs/neighbors/ball_cover-ext.cuh b/cpp/include/cuvs/neighbors/ball_cover-ext.cuh deleted file mode 100644 index b1cd2b4ed..000000000 --- a/cpp/include/cuvs/neighbors/ball_cover-ext.cuh +++ /dev/null @@ -1,124 +0,0 @@ -/* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#include // uint32_t -#include // cuvs::distance::DistanceType -#include // BallCoverIndex -#include // RAFT_EXPLICIT - -#ifdef RAFT_EXPLICIT_INSTANTIATE_ONLY - -namespace cuvs::neighbors::ball_cover { - -template -void build_index(raft::resources const& handle, - BallCoverIndex& index) RAFT_EXPLICIT; - -template -void all_knn_query(raft::resources const& handle, - BallCoverIndex& index, - int_t k, - idx_t* inds, - value_t* dists, - bool perform_post_filtering = true, - float weight = 1.0) RAFT_EXPLICIT; - -template -void all_knn_query(raft::resources const& handle, - BallCoverIndex& index, - raft::device_matrix_view inds, - raft::device_matrix_view dists, - int_t k, - bool perform_post_filtering = true, - float weight = 1.0) RAFT_EXPLICIT; - -template -void knn_query(raft::resources const& handle, - const BallCoverIndex& index, - int_t k, - const value_t* query, - int_t n_query_pts, - idx_t* inds, - value_t* dists, - bool perform_post_filtering = true, - float weight = 1.0) RAFT_EXPLICIT; - -template -void knn_query(raft::resources const& handle, - const BallCoverIndex& index, - raft::device_matrix_view query, - raft::device_matrix_view inds, - raft::device_matrix_view dists, - int_t k, - bool perform_post_filtering = true, - float weight = 1.0) RAFT_EXPLICIT; - -} // namespace cuvs::neighbors::ball_cover - -#endif // RAFT_EXPLICIT_INSTANTIATE_ONLY - -#define instantiate_raft_neighbors_ball_cover(idx_t, value_t, int_t, matrix_idx_t) \ - extern template void \ - cuvs::neighbors::ball_cover::build_index( \ - raft::resources const& handle, \ - cuvs::neighbors::ball_cover::BallCoverIndex& index); \ - \ - extern template void \ - cuvs::neighbors::ball_cover::all_knn_query( \ - raft::resources const& handle, \ - cuvs::neighbors::ball_cover::BallCoverIndex& index, \ - int_t k, \ - idx_t* inds, \ - value_t* dists, \ - bool perform_post_filtering, \ - float weight); \ - \ - extern template void \ - cuvs::neighbors::ball_cover::all_knn_query( \ - raft::resources const& handle, \ - cuvs::neighbors::ball_cover::BallCoverIndex& index, \ - raft::device_matrix_view inds, \ - raft::device_matrix_view dists, \ - int_t k, \ - bool perform_post_filtering, \ - float weight); \ - \ - extern template void cuvs::neighbors::ball_cover::knn_query( \ - raft::resources const& handle, \ - const cuvs::neighbors::ball_cover::BallCoverIndex& index, \ - int_t k, \ - const value_t* query, \ - int_t n_query_pts, \ - idx_t* inds, \ - value_t* dists, \ - bool perform_post_filtering, \ - float weight); \ - \ - extern template void \ - cuvs::neighbors::ball_cover::knn_query( \ - raft::resources const& handle, \ - const cuvs::neighbors::ball_cover::BallCoverIndex& index, \ - raft::device_matrix_view query, \ - raft::device_matrix_view inds, \ - raft::device_matrix_view dists, \ - int_t k, \ - bool perform_post_filtering, \ - float weight); - -instantiate_raft_neighbors_ball_cover(int64_t, float, uint32_t, uint32_t); - -#undef instantiate_raft_neighbors_ball_cover diff --git a/cpp/include/cuvs/neighbors/ball_cover-inl.cuh b/cpp/include/cuvs/neighbors/ball_cover-inl.cuh deleted file mode 100644 index 4d0f170df..000000000 --- a/cpp/include/cuvs/neighbors/ball_cover-inl.cuh +++ /dev/null @@ -1,395 +0,0 @@ -/* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#ifndef __BALL_COVER_H -#define __BALL_COVER_H - -#pragma once - -#include - -#include -#include -#include -#include -#include - -namespace cuvs::neighbors::ball_cover { - -/** - * @defgroup random_ball_cover Random Ball Cover algorithm - * @{ - */ - -/** - * Builds and populates a previously unbuilt BallCoverIndex - * - * Usage example: - * @code{.cpp} - * - * #include - * #include - * #include - * using namespace cuvs::neighbors; - * - * raft::resources handle; - * ... - * auto metric = cuvs::distance::DistanceType::L2Expanded; - * BallCoverIndex index(handle, X, metric); - * - * ball_cover::build_index(handle, index); - * @endcode - * - * @tparam idx_t knn index type - * @tparam value_t knn value type - * @tparam int_t integral type for knn params - * @tparam matrix_idx_t matrix indexing type - * @param[in] handle library resource management handle - * @param[inout] index an empty (and not previous built) instance of BallCoverIndex - */ -template -void build_index(raft::resources const& handle, - BallCoverIndex& index) -{ - ASSERT(index.n <= 3, "only 2d and 3d vectors are supported in current implementation"); - if (index.metric == cuvs::distance::DistanceType::Haversine) { - cuvs::spatial::knn::detail::rbc_build_index( - handle, index, spatial::knn::detail::HaversineFunc()); - } else if (index.metric == cuvs::distance::DistanceType::L2SqrtExpanded || - index.metric == cuvs::distance::DistanceType::L2SqrtUnexpanded) { - cuvs::spatial::knn::detail::rbc_build_index( - handle, index, spatial::knn::detail::EuclideanFunc()); - } else { - RAFT_FAIL("Metric not support"); - } - - index.set_index_trained(); -} - -/** @} */ // end group random_ball_cover - -/** - * Performs a faster exact knn in metric spaces using the triangle - * inequality with a number of landmark points to reduce the - * number of distance computations from O(n^2) to O(sqrt(n)). This - * performs an all neighbors knn, which can reuse memory when - * the index and query are the same array. This function will - * build the index and assumes rbc_build_index() has not already - * been called. - * @tparam idx_t knn index type - * @tparam value_t knn distance type - * @tparam int_t type for integers, such as number of rows/cols - * @param[in] handle raft handle for resource management - * @param[inout] index ball cover index which has not yet been built - * @param[in] k number of nearest neighbors to find - * @param[in] perform_post_filtering if this is false, only the closest k landmarks - * are considered (which will return approximate - * results). - * @param[out] inds output knn indices - * @param[out] dists output knn distances - * @param[in] weight a weight for overlap between the closest landmark and - * the radius of other landmarks when pruning distances. - * Setting this value below 1 can effectively turn off - * computing distances against many other balls, enabling - * approximate nearest neighbors. Recall can be adjusted - * based on how many relevant balls are ignored. Note that - * many datasets can still have great recall even by only - * looking in the closest landmark. - */ -template -void all_knn_query(raft::resources const& handle, - BallCoverIndex& index, - int_t k, - idx_t* inds, - value_t* dists, - bool perform_post_filtering = true, - float weight = 1.0) -{ - ASSERT(index.n <= 3, "only 2d and 3d vectors are supported in current implementation"); - if (index.metric == cuvs::distance::DistanceType::Haversine) { - cuvs::spatial::knn::detail::rbc_all_knn_query( - handle, - index, - k, - inds, - dists, - spatial::knn::detail::HaversineFunc(), - perform_post_filtering, - weight); - } else if (index.metric == cuvs::distance::DistanceType::L2SqrtExpanded || - index.metric == cuvs::distance::DistanceType::L2SqrtUnexpanded) { - cuvs::spatial::knn::detail::rbc_all_knn_query( - handle, - index, - k, - inds, - dists, - spatial::knn::detail::EuclideanFunc(), - perform_post_filtering, - weight); - } else { - RAFT_FAIL("Metric not supported"); - } - - index.set_index_trained(); -} - -/** - * @ingroup random_ball_cover - * @{ - */ - -/** - * Performs a faster exact knn in metric spaces using the triangle - * inequality with a number of landmark points to reduce the - * number of distance computations from O(n^2) to O(sqrt(n)). This - * performs an all neighbors knn, which can reuse memory when - * the index and query are the same array. This function will - * build the index and assumes rbc_build_index() has not already - * been called. - * - * Usage example: - * @code{.cpp} - * - * #include - * #include - * #include - * using namespace cuvs::neighbors; - * - * raft::resources handle; - * ... - * auto metric = cuvs::distance::DistanceType::L2Expanded; - * - * // Construct a ball cover index - * BallCoverIndex index(handle, X, metric); - * - * // Perform all neighbors knn query - * ball_cover::all_knn_query(handle, index, inds, dists, k); - * @endcode - * - * @tparam idx_t knn index type - * @tparam value_t knn distance type - * @tparam int_t type for integers, such as number of rows/cols - * @tparam matrix_idx_t matrix indexing type - * - * @param[in] handle raft handle for resource management - * @param[in] index ball cover index which has not yet been built - * @param[out] inds output knn indices - * @param[out] dists output knn distances - * @param[in] k number of nearest neighbors to find - * @param[in] perform_post_filtering if this is false, only the closest k landmarks - * are considered (which will return approximate - * results). - * @param[in] weight a weight for overlap between the closest landmark and - * the radius of other landmarks when pruning distances. - * Setting this value below 1 can effectively turn off - * computing distances against many other balls, enabling - * approximate nearest neighbors. Recall can be adjusted - * based on how many relevant balls are ignored. Note that - * many datasets can still have great recall even by only - * looking in the closest landmark. - */ -template -void all_knn_query(raft::resources const& handle, - BallCoverIndex& index, - raft::device_matrix_view inds, - raft::device_matrix_view dists, - int_t k, - bool perform_post_filtering = true, - float weight = 1.0) -{ - RAFT_EXPECTS(index.n <= 3, "only 2d and 3d vectors are supported in current implementation"); - RAFT_EXPECTS(k <= index.m, - "k must be less than or equal to the number of data points in the index"); - RAFT_EXPECTS(inds.extent(1) == dists.extent(1) && dists.extent(1) == static_cast(k), - "Number of columns in output indices and distances matrices must be equal to k"); - - RAFT_EXPECTS(inds.extent(0) == dists.extent(0) && dists.extent(0) == index.get_X().extent(0), - "Number of rows in output indices and distances matrices must equal number of rows " - "in index matrix."); - - all_knn_query( - handle, index, k, inds.data_handle(), dists.data_handle(), perform_post_filtering, weight); -} - -/** @} */ - -/** - * Performs a faster exact knn in metric spaces using the triangle - * inequality with a number of landmark points to reduce the - * number of distance computations from O(n^2) to O(sqrt(n)). This - * function does not build the index and assumes rbc_build_index() has - * already been called. Use this function when the index and - * query arrays are different, otherwise use rbc_all_knn_query(). - * @tparam idx_t index type - * @tparam value_t distances type - * @tparam int_t integer type for size info - * @param[in] handle raft handle for resource management - * @param[inout] index ball cover index which has not yet been built - * @param[in] k number of nearest neighbors to find - * @param[in] query the - * @param[in] perform_post_filtering if this is false, only the closest k landmarks - * are considered (which will return approximate - * results). - * @param[out] inds output knn indices - * @param[out] dists output knn distances - * @param[in] weight a weight for overlap between the closest landmark and - * the radius of other landmarks when pruning distances. - * Setting this value below 1 can effectively turn off - * computing distances against many other balls, enabling - * approximate nearest neighbors. Recall can be adjusted - * based on how many relevant balls are ignored. Note that - * many datasets can still have great recall even by only - * looking in the closest landmark. - * @param[in] n_query_pts number of query points - */ -template -void knn_query(raft::resources const& handle, - const BallCoverIndex& index, - int_t k, - const value_t* query, - int_t n_query_pts, - idx_t* inds, - value_t* dists, - bool perform_post_filtering = true, - float weight = 1.0) -{ - ASSERT(index.n <= 3, "only 2d and 3d vectors are supported in current implementation"); - if (index.metric == cuvs::distance::DistanceType::Haversine) { - cuvs::spatial::knn::detail::rbc_knn_query(handle, - index, - k, - query, - n_query_pts, - inds, - dists, - spatial::knn::detail::HaversineFunc(), - perform_post_filtering, - weight); - } else if (index.metric == cuvs::distance::DistanceType::L2SqrtExpanded || - index.metric == cuvs::distance::DistanceType::L2SqrtUnexpanded) { - cuvs::spatial::knn::detail::rbc_knn_query(handle, - index, - k, - query, - n_query_pts, - inds, - dists, - spatial::knn::detail::EuclideanFunc(), - perform_post_filtering, - weight); - } else { - RAFT_FAIL("Metric not supported"); - } -} - -/** - * @ingroup random_ball_cover - * @{ - */ - -/** - * Performs a faster exact knn in metric spaces using the triangle - * inequality with a number of landmark points to reduce the - * number of distance computations from O(n^2) to O(sqrt(n)). This - * function does not build the index and assumes rbc_build_index() has - * already been called. Use this function when the index and - * query arrays are different, otherwise use rbc_all_knn_query(). - * - * Usage example: - * @code{.cpp} - * - * #include - * #include - * #include - * using namespace cuvs::neighbors; - * - * raft::resources handle; - * ... - * auto metric = cuvs::distance::DistanceType::L2Expanded; - * - * // Build a ball cover index - * BallCoverIndex index(handle, X, metric); - * ball_cover::build_index(handle, index); - * - * // Perform all neighbors knn query - * ball_cover::knn_query(handle, index, inds, dists, k); - * @endcode - - * - * @tparam idx_t index type - * @tparam value_t distances type - * @tparam int_t integer type for size info - * @tparam matrix_idx_t - * @param[in] handle raft handle for resource management - * @param[in] index ball cover index which has not yet been built - * @param[in] query device matrix containing query data points - * @param[out] inds output knn indices - * @param[out] dists output knn distances - * @param[in] k number of nearest neighbors to find - * @param[in] perform_post_filtering if this is false, only the closest k landmarks - * are considered (which will return approximate - * results). - * @param[in] weight a weight for overlap between the closest landmark and - * the radius of other landmarks when pruning distances. - * Setting this value below 1 can effectively turn off - * computing distances against many other balls, enabling - * approximate nearest neighbors. Recall can be adjusted - * based on how many relevant balls are ignored. Note that - * many datasets can still have great recall even by only - * looking in the closest landmark. - */ -template -void knn_query(raft::resources const& handle, - const BallCoverIndex& index, - raft::device_matrix_view query, - raft::device_matrix_view inds, - raft::device_matrix_view dists, - int_t k, - bool perform_post_filtering = true, - float weight = 1.0) -{ - RAFT_EXPECTS(k <= index.m, - "k must be less than or equal to the number of data points in the index"); - RAFT_EXPECTS(inds.extent(1) == dists.extent(1) && dists.extent(1) == static_cast(k), - "Number of columns in output indices and distances matrices must be equal to k"); - - RAFT_EXPECTS(inds.extent(0) == dists.extent(0) && dists.extent(0) == query.extent(0), - "Number of rows in output indices and distances matrices must equal number of rows " - "in search matrix."); - - RAFT_EXPECTS(query.extent(1) == index.get_X().extent(1), - "Number of columns in query and index matrices must match."); - - knn_query(handle, - index, - k, - query.data_handle(), - query.extent(0), - inds.data_handle(), - dists.data_handle(), - perform_post_filtering, - weight); -} - -/** @} */ - -// TODO: implement functions for: -// 4. rbc_eps_neigh() - given a populated index, perform query against different query array -// 5. rbc_all_eps_neigh() - populate a BallCoverIndex and query against training data - -} // namespace cuvs::neighbors::ball_cover - -#endif diff --git a/cpp/include/cuvs/neighbors/ball_cover.cuh b/cpp/include/cuvs/neighbors/ball_cover.cuh deleted file mode 100644 index 41c5d0310..000000000 --- a/cpp/include/cuvs/neighbors/ball_cover.cuh +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#ifndef RAFT_EXPLICIT_INSTANTIATE_ONLY -#include "ball_cover-inl.cuh" -#endif - -#ifdef RAFT_COMPILED -#include "ball_cover-ext.cuh" -#endif diff --git a/cpp/include/cuvs/neighbors/ball_cover_types.hpp b/cpp/include/cuvs/neighbors/ball_cover_types.hpp deleted file mode 100644 index c6e9fab2c..000000000 --- a/cpp/include/cuvs/neighbors/ball_cover_types.hpp +++ /dev/null @@ -1,169 +0,0 @@ -/* - * Copyright (c) 2021-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include -#include -#include -#include -#include -#include - -namespace cuvs::neighbors::ball_cover { - -/** - * @ingroup random_ball_cover - * @{ - */ - -/** - * Stores raw index data points, sampled landmarks, the 1-nns of index points - * to their closest landmarks, and the ball radii of each landmark. This - * class is intended to be constructed once and reused across subsequent - * queries. - * @tparam value_idx - * @tparam value_t - * @tparam value_int - */ -template -class BallCoverIndex { - public: - explicit BallCoverIndex(raft::resources const& handle_, - const value_t* X_, - value_int m_, - value_int n_, - cuvs::distance::DistanceType metric_) - : handle(handle_), - X(raft::make_device_matrix_view(X_, m_, n_)), - m(m_), - n(n_), - metric(metric_), - /** - * the sqrt() here makes the sqrt(m)^2 a linear-time lower bound - * - * Total memory footprint of index: (2 * sqrt(m)) + (n * sqrt(m)) + (2 * m) - */ - n_landmarks(sqrt(m_)), - R_indptr(raft::make_device_vector(handle, sqrt(m_) + 1)), - R_1nn_cols(raft::make_device_vector(handle, m_)), - R_1nn_dists(raft::make_device_vector(handle, m_)), - R_closest_landmark_dists(raft::make_device_vector(handle, m_)), - R(raft::make_device_matrix(handle, sqrt(m_), n_)), - R_radius(raft::make_device_vector(handle, sqrt(m_))), - index_trained(false) - { - } - - explicit BallCoverIndex(raft::resources const& handle_, - raft::device_matrix_view X_, - cuvs::distance::DistanceType metric_) - : handle(handle_), - X(X_), - m(X_.extent(0)), - n(X_.extent(1)), - metric(metric_), - /** - * the sqrt() here makes the sqrt(m)^2 a linear-time lower bound - * - * Total memory footprint of index: (2 * sqrt(m)) + (n * sqrt(m)) + (2 * m) - */ - n_landmarks(sqrt(X_.extent(0))), - R_indptr(raft::make_device_vector(handle, sqrt(X_.extent(0)) + 1)), - R_1nn_cols(raft::make_device_vector(handle, X_.extent(0))), - R_1nn_dists(raft::make_device_vector(handle, X_.extent(0))), - R_closest_landmark_dists(raft::make_device_vector(handle, X_.extent(0))), - R(raft::make_device_matrix(handle, sqrt(X_.extent(0)), X_.extent(1))), - R_radius(raft::make_device_vector(handle, sqrt(X_.extent(0)))), - index_trained(false) - { - } - - auto get_R_indptr() const -> raft::device_vector_view - { - return R_indptr.view(); - } - auto get_R_1nn_cols() const -> raft::device_vector_view - { - return R_1nn_cols.view(); - } - auto get_R_1nn_dists() const -> raft::device_vector_view - { - return R_1nn_dists.view(); - } - auto get_R_radius() const -> raft::device_vector_view - { - return R_radius.view(); - } - auto get_R() const -> raft::device_matrix_view - { - return R.view(); - } - auto get_R_closest_landmark_dists() const -> raft::device_vector_view - { - return R_closest_landmark_dists.view(); - } - - raft::device_vector_view get_R_indptr() { return R_indptr.view(); } - raft::device_vector_view get_R_1nn_cols() { return R_1nn_cols.view(); } - raft::device_vector_view get_R_1nn_dists() { return R_1nn_dists.view(); } - raft::device_vector_view get_R_radius() { return R_radius.view(); } - raft::device_matrix_view get_R() { return R.view(); } - raft::device_vector_view get_R_closest_landmark_dists() - { - return R_closest_landmark_dists.view(); - } - raft::device_matrix_view get_X() const { return X; } - - cuvs::distance::DistanceType get_metric() const { return metric; } - - value_int get_n_landmarks() const { return n_landmarks; } - bool is_index_trained() const { return index_trained; }; - - // This should only be set by internal functions - void set_index_trained() { index_trained = true; } - - raft::resources const& handle; - - value_int m; - value_int n; - value_int n_landmarks; - - raft::device_matrix_view X; - - cuvs::distance::DistanceType metric; - - private: - // CSR storing the neighborhoods for each data point - raft::device_vector R_indptr; - raft::device_vector R_1nn_cols; - raft::device_vector R_1nn_dists; - raft::device_vector R_closest_landmark_dists; - - raft::device_vector R_radius; - - raft::device_matrix R; - - protected: - bool index_trained; -}; - -/** @} */ - -} // namespace cuvs::neighbors::ball_cover diff --git a/cpp/include/cuvs/neighbors/brute_force-ext.cuh b/cpp/include/cuvs/neighbors/brute_force-ext.cuh deleted file mode 100644 index bc4773513..000000000 --- a/cpp/include/cuvs/neighbors/brute_force-ext.cuh +++ /dev/null @@ -1,149 +0,0 @@ -/* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include - -#include // cuvs::distance::DistanceType -#include -#include // raft::device_matrix_view -#include // raft::identity_op -#include // raft::resources -#include // RAFT_EXPLICIT - -#ifdef RAFT_EXPLICIT_INSTANTIATE_ONLY - -namespace cuvs::neighbors::brute_force { - -template -inline void knn_merge_parts( - raft::resources const& handle, - raft::device_matrix_view in_keys, - raft::device_matrix_view in_values, - raft::device_matrix_view out_keys, - raft::device_matrix_view out_values, - size_t n_samples, - std::optional> translations = std::nullopt) RAFT_EXPLICIT; - -template -index build( - raft::resources const& res, - raft::mdspan, raft::row_major, Accessor> dataset, - cuvs::distance::DistanceType metric = distance::DistanceType::L2Unexpanded, - T metric_arg = 0.0) RAFT_EXPLICIT; - -template -void search(raft::resources const& res, - const index& idx, - raft::device_matrix_view queries, - raft::device_matrix_view neighbors, - raft::device_matrix_view distances) RAFT_EXPLICIT; - -template -void knn(raft::resources const& handle, - std::vector> index, - raft::device_matrix_view search, - raft::device_matrix_view indices, - raft::device_matrix_view distances, - distance::DistanceType metric = distance::DistanceType::L2Unexpanded, - std::optional metric_arg = std::make_optional(2.0f), - std::optional global_id_offset = std::nullopt, - epilogue_op distance_epilogue = raft::identity_op()) RAFT_EXPLICIT; - -template -void fused_l2_knn(raft::resources const& handle, - raft::device_matrix_view index, - raft::device_matrix_view query, - raft::device_matrix_view out_inds, - raft::device_matrix_view out_dists, - cuvs::distance::DistanceType metric) RAFT_EXPLICIT; - -} // namespace cuvs::neighbors::brute_force - -#endif // RAFT_EXPLICIT_INSTANTIATE_ONLY - -// No extern template for cuvs::neighbors::brute_force::knn_merge_parts - -#define instantiate_raft_neighbors_brute_force_knn( \ - idx_t, value_t, matrix_idx, index_layout, search_layout, epilogue_op) \ - extern template void cuvs::neighbors::brute_force:: \ - knn( \ - raft::resources const& handle, \ - std::vector> index, \ - raft::device_matrix_view search, \ - raft::device_matrix_view indices, \ - raft::device_matrix_view distances, \ - cuvs::distance::DistanceType metric, \ - std::optional metric_arg, \ - std::optional global_id_offset, \ - epilogue_op distance_epilogue); - -instantiate_raft_neighbors_brute_force_knn( - int64_t, float, uint32_t, raft::row_major, raft::row_major, raft::identity_op); -instantiate_raft_neighbors_brute_force_knn( - int64_t, float, int64_t, raft::row_major, raft::row_major, raft::identity_op); -instantiate_raft_neighbors_brute_force_knn( - int, float, int, raft::row_major, raft::row_major, raft::identity_op); -instantiate_raft_neighbors_brute_force_knn( - uint32_t, float, uint32_t, raft::row_major, raft::row_major, raft::identity_op); - -#undef instantiate_raft_neighbors_brute_force_knn - -namespace cuvs::neighbors::brute_force { - -extern template void search( - raft::resources const& res, - const cuvs::neighbors::brute_force::index& idx, - raft::device_matrix_view queries, - raft::device_matrix_view neighbors, - raft::device_matrix_view distances); - -extern template void search( - raft::resources const& res, - const cuvs::neighbors::brute_force::index& idx, - raft::device_matrix_view queries, - raft::device_matrix_view neighbors, - raft::device_matrix_view distances); - -extern template cuvs::neighbors::brute_force::index build( - raft::resources const& res, - raft::device_matrix_view dataset, - cuvs::distance::DistanceType metric, - float metric_arg); -} // namespace cuvs::neighbors::brute_force - -#define instantiate_raft_neighbors_brute_force_fused_l2_knn( \ - value_t, idx_t, idx_layout, query_layout) \ - extern template void cuvs::neighbors::brute_force::fused_l2_knn( \ - raft::resources const& handle, \ - raft::device_matrix_view index, \ - raft::device_matrix_view query, \ - raft::device_matrix_view out_inds, \ - raft::device_matrix_view out_dists, \ - cuvs::distance::DistanceType metric); - -instantiate_raft_neighbors_brute_force_fused_l2_knn(float, - int64_t, - raft::row_major, - raft::row_major) - -#undef instantiate_raft_neighbors_brute_force_fused_l2_knn diff --git a/cpp/include/cuvs/neighbors/brute_force-inl.cuh b/cpp/include/cuvs/neighbors/brute_force-inl.cuh deleted file mode 100644 index 3d5c449a9..000000000 --- a/cpp/include/cuvs/neighbors/brute_force-inl.cuh +++ /dev/null @@ -1,355 +0,0 @@ -/* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include -#include -#include -#include -#include -#include - -namespace cuvs::neighbors::brute_force { - -/** - * @defgroup brute_force_knn Brute-force K-Nearest Neighbors - * @{ - */ - -/** - * @brief Performs a k-select across several (contiguous) row-partitioned index/distance - * matrices formatted like the following: - * - * part1row1: k0, k1, k2, k3 - * part1row2: k0, k1, k2, k3 - * part1row3: k0, k1, k2, k3 - * part2row1: k0, k1, k2, k3 - * part2row2: k0, k1, k2, k3 - * part2row3: k0, k1, k2, k3 - * etc... - * - * The example above shows what an aggregated index/distance matrix - * would look like with two partitions when n_samples=3 and k=4. - * - * When working with extremely large data sets that have been broken - * over multiple indexes, such as when computing over multiple GPUs, - * the ids will often start at 0 for each local knn index but the - * global ids need to be used when merging them together. An optional - * translations vector can be supplied to map the starting id of - * each partition to its global id so that the final merged knn - * is based on the global ids. - * - * Usage example: - * @code{.cpp} - * #include - * #include - * using namespace cuvs::neighbors; - * - * raft::resources handle; - * ... - * compute multiple knn graphs and aggregate row-wise - * (see detailed description above) - * ... - * brute_force::knn_merge_parts(handle, in_keys, in_values, out_keys, out_values, n_samples); - * @endcode - * - * @tparam idx_t - * @tparam value_t - * - * @param[in] handle - * @param[in] in_keys matrix of input keys (size n_samples * n_parts * k) - * @param[in] in_values matrix of input values (size n_samples * n_parts * k) - * @param[out] out_keys matrix of output keys (size n_samples * k) - * @param[out] out_values matrix of output values (size n_samples * k) - * @param[in] n_samples number of rows in each partition - * @param[in] translations optional vector of starting global id mappings for each local partition - */ -template -inline void knn_merge_parts( - raft::resources const& handle, - raft::device_matrix_view in_keys, - raft::device_matrix_view in_values, - raft::device_matrix_view out_keys, - raft::device_matrix_view out_values, - size_t n_samples, - std::optional> translations = std::nullopt) -{ - RAFT_EXPECTS(in_keys.extent(1) == in_values.extent(1) && in_keys.extent(0) == in_values.extent(0), - "in_keys and in_values must have the same shape."); - RAFT_EXPECTS( - out_keys.extent(0) == out_values.extent(0) && out_keys.extent(0) == n_samples, - "Number of rows in output keys and val matrices must equal number of rows in search matrix."); - RAFT_EXPECTS( - out_keys.extent(1) == out_values.extent(1) && out_keys.extent(1) == in_keys.extent(1), - "Number of columns in output indices and distances matrices must be equal to k"); - - idx_t* translations_ptr = nullptr; - if (translations.has_value()) { translations_ptr = translations.value().data_handle(); } - - auto n_parts = in_keys.extent(0) / n_samples; - detail::knn_merge_parts(in_keys.data_handle(), - in_values.data_handle(), - out_keys.data_handle(), - out_values.data_handle(), - n_samples, - n_parts, - in_keys.extent(1), - resource::get_cuda_stream(handle), - translations_ptr); -} - -/** - * @brief Flat C++ API function to perform a brute force knn on - * a series of input arrays and combine the results into a single - * output array for indexes and distances. Inputs can be either - * row- or column-major but the output matrices will always be in - * row-major format. - * - * Usage example: - * @code{.cpp} - * #include - * #include - * #include - * using namespace cuvs::neighbors; - * - * raft::resources handle; - * ... - * auto metric = cuvs::distance::DistanceType::L2SqrtExpanded; - * brute_force::knn(handle, index, search, indices, distances, metric); - * @endcode - * - * @param[in] handle: the cuml handle to use - * @param[in] index: vector of device matrices (each size m_i*d) to be used as the knn index - * @param[in] search: matrix (size n*d) to be used for searching the index - * @param[out] indices: matrix (size n*k) to store output knn indices - * @param[out] distances: matrix (size n*k) to store the output knn distance - * @param[in] metric: distance metric to use. Euclidean (L2) is used by default - * @param[in] metric_arg: the value of `p` for Minkowski (l-p) distances. This - * is ignored if the metric_type is not Minkowski. - * @param[in] global_id_offset: optional starting global id mapping for the local partition - * (assumes the index contains contiguous ids in the global id space) - * @param[in] distance_epilogue: optional epilogue function to run after computing distances. This - function takes a triple of the (value, rowid, colid) for each - element in the pairwise distances and returns a transformed value - back. - */ -template -void knn(raft::resources const& handle, - std::vector> index, - raft::device_matrix_view search, - raft::device_matrix_view indices, - raft::device_matrix_view distances, - distance::DistanceType metric = distance::DistanceType::L2Unexpanded, - std::optional metric_arg = std::make_optional(2.0f), - std::optional global_id_offset = std::nullopt, - epilogue_op distance_epilogue = raft::identity_op()) -{ - RAFT_EXPECTS(index[0].extent(1) == search.extent(1), - "Number of dimensions for both index and search matrices must be equal"); - - RAFT_EXPECTS(indices.extent(0) == distances.extent(0) && distances.extent(0) == search.extent(0), - "Number of rows in output indices and distances matrices must equal number of rows " - "in search matrix."); - RAFT_EXPECTS(indices.extent(1) == distances.extent(1) && distances.extent(1), - "Number of columns in output indices and distances matrices must the same"); - - bool rowMajorIndex = std::is_same_v; - bool rowMajorQuery = std::is_same_v; - - std::vector inputs; - std::vector sizes; - for (std::size_t i = 0; i < index.size(); ++i) { - inputs.push_back(const_cast(index[i].data_handle())); - sizes.push_back(index[i].extent(0)); - } - - std::vector trans; - if (global_id_offset.has_value()) { trans.push_back(global_id_offset.value()); } - - std::vector* trans_arg = global_id_offset.has_value() ? &trans : nullptr; - - cuvs::neighbors::detail::brute_force_knn_impl(handle, - inputs, - sizes, - index[0].extent(1), - // TODO: This is unfortunate. Need to fix. - const_cast(search.data_handle()), - search.extent(0), - indices.data_handle(), - distances.data_handle(), - indices.extent(1), - rowMajorIndex, - rowMajorQuery, - trans_arg, - metric, - metric_arg.value_or(2.0f), - distance_epilogue); -} - -/** - * @brief Compute the k-nearest neighbors using L2 expanded/unexpanded distance. - * - * This is a specialized function for fusing the k-selection with the distance - * computation when k < 64. The value of k will be inferred from the number - * of columns in the output matrices. - * - * Usage example: - * @code{.cpp} - * #include - * #include - * #include - * using namespace cuvs::neighbors; - * - * raft::resources handle; - * ... - * auto metric = cuvs::distance::DistanceType::L2SqrtExpanded; - * brute_force::fused_l2_knn(handle, index, search, indices, distances, metric); - * @endcode - - * @tparam value_t type of values - * @tparam idx_t type of indices - * @tparam idx_layout layout type of index matrix - * @tparam query_layout layout type of query matrix - * @param[in] handle raft handle for sharing expensive resources - * @param[in] index input index array on device (size m * d) - * @param[in] query input query array on device (size n * d) - * @param[out] out_inds output indices array on device (size n * k) - * @param[out] out_dists output dists array on device (size n * k) - * @param[in] metric type of distance computation to perform (must be a variant of L2) - */ -template -void fused_l2_knn(raft::resources const& handle, - raft::device_matrix_view index, - raft::device_matrix_view query, - raft::device_matrix_view out_inds, - raft::device_matrix_view out_dists, - cuvs::distance::DistanceType metric) -{ - int k = static_cast(out_inds.extent(1)); - - RAFT_EXPECTS(k <= 64, "For fused k-selection, k must be < 64"); - RAFT_EXPECTS(out_inds.extent(1) == out_dists.extent(1), "Value of k must match for outputs"); - RAFT_EXPECTS(index.extent(1) == query.extent(1), - "Number of columns in input matrices must be the same."); - - RAFT_EXPECTS(metric == distance::DistanceType::L2Expanded || - metric == distance::DistanceType::L2Unexpanded || - metric == distance::DistanceType::L2SqrtUnexpanded || - metric == distance::DistanceType::L2SqrtExpanded, - "Distance metric must be L2"); - - size_t n_index_rows = index.extent(0); - size_t n_query_rows = query.extent(0); - size_t D = index.extent(1); - - RAFT_EXPECTS(raft::is_row_or_column_major(index), "Index must be row or column major layout"); - RAFT_EXPECTS(raft::is_row_or_column_major(query), "Query must be row or column major layout"); - - const bool rowMajorIndex = raft::is_row_major(index); - const bool rowMajorQuery = raft::is_row_major(query); - - cuvs::spatial::knn::detail::fusedL2Knn(D, - out_inds.data_handle(), - out_dists.data_handle(), - index.data_handle(), - query.data_handle(), - n_index_rows, - n_query_rows, - k, - rowMajorIndex, - rowMajorQuery, - raft::resource::get_cuda_stream(handle), - metric); -} - -/** - * @brief Build the index from the dataset for efficient search. - * - * @tparam T data element type - * - * @param[in] res - * @param[in] dataset a matrix view (host or device) to a row-major matrix [n_rows, dim] - * @param[in] metric: distance metric to use. Euclidean (L2) is used by default - * @param[in] metric_arg: the value of `p` for Minkowski (l-p) distances. This - * is ignored if the metric_type is not Minkowski. - * - * @return the constructed brute force index - */ -template -index build( - raft::resources const& res, - raft::mdspan, raft::row_major, Accessor> dataset, - cuvs::distance::DistanceType metric = distance::DistanceType::L2Unexpanded, - T metric_arg = 0.0) -{ - // certain distance metrics can benefit by pre-calculating the norms for the index dataset - // which lets us avoid calculating these at query time - std::optional> norms; - if (metric == cuvs::distance::DistanceType::L2Expanded || - metric == cuvs::distance::DistanceType::L2SqrtExpanded || - metric == cuvs::distance::DistanceType::CosineExpanded) { - norms = raft::make_device_vector(res, dataset.extent(0)); - // cosine needs the l2norm, where as l2 distances needs the squared norm - if (metric == cuvs::distance::DistanceType::CosineExpanded) { - raft::linalg::norm(res, - dataset, - norms->view(), - raft::linalg::NormType::L2Norm, - raft::linalg::Apply::ALONG_ROWS, - raft::sqrt_op{}); - } else { - raft::linalg::norm(res, - dataset, - norms->view(), - raft::linalg::NormType::L2Norm, - raft::linalg::Apply::ALONG_ROWS); - } - } - - return index(res, dataset, std::move(norms), metric, metric_arg); -} - -/** - * @brief Brute Force search using the constructed index. - * - * @tparam T data element type - * @tparam IdxT type of the indices - * - * @param[in] res raft resources - * @param[in] idx brute force index - * @param[in] queries a device matrix view to a row-major matrix [n_queries, index->dim()] - * @param[out] neighbors a device matrix view to the indices of the neighbors in the source dataset - * [n_queries, k] - * @param[out] distances a device matrix view to the distances to the selected neighbors [n_queries, - * k] - */ -template -void search(raft::resources const& res, - const index& idx, - raft::device_matrix_view queries, - raft::device_matrix_view neighbors, - raft::device_matrix_view distances) -{ - cuvs::neighbors::detail::brute_force_search(res, idx, queries, neighbors, distances); -} -/** @} */ // end group brute_force_knn -} // namespace cuvs::neighbors::brute_force diff --git a/cpp/include/cuvs/neighbors/brute_force.cuh b/cpp/include/cuvs/neighbors/brute_force.cuh deleted file mode 100644 index 91065d35f..000000000 --- a/cpp/include/cuvs/neighbors/brute_force.cuh +++ /dev/null @@ -1,92 +0,0 @@ -/* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once -#include - -#ifndef RAFT_EXPLICIT_INSTANTIATE_ONLY -#include "brute_force-inl.cuh" -#endif - -#ifdef RAFT_COMPILED -#include "brute_force-ext.cuh" -#endif - -#include - -namespace cuvs::neighbors::brute_force { -/** - * @brief Make a brute force query over batches of k - * - * This lets you query for batches of k. For example, you can get - * the first 100 neighbors, then the next 100 neighbors etc. - * - * Example usage: - * @code{.cpp} - * #include - * #include - * #include - - * // create a random dataset - * int n_rows = 10000; - * int n_cols = 10000; - - * raft::device_resources res; - * auto dataset = raft::make_device_matrix(res, n_rows, n_cols); - * auto labels = raft::make_device_vector(res, n_rows); - - * raft::make_blobs(res, dataset.view(), labels.view()); - * - * // create a brute_force knn index from the dataset - * auto index = cuvs::neighbors::brute_force::build(res, - * raft::make_const_mdspan(dataset.view())); - * - * // search the index in batches of 128 nearest neighbors - * auto search = raft::make_const_mdspan(dataset.view()); - * auto query = make_batch_k_query(res, index, search, 128); - * for (auto & batch: *query) { - * // batch.indices() and batch.distances() contain the information on the current batch - * } - * - * // we can also support variable sized batches - loaded up a different number - * // of neighbors at each iteration through the ::advance method - * int64_t batch_size = 128; - * query = make_batch_k_query(res, index, search, batch_size); - * for (auto it = query->begin(); it != query->end(); it.advance(batch_size)) { - * // batch.indices() and batch.distances() contain the information on the current batch - * - * batch_size += 16; // load up an extra 16 items in the next batch - * } - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset - * @param[in] res - * @param[in] index The index to query - * @param[in] query A device matrix view to query for [n_queries, index->dim()] - * @param[in] batch_size The size of each batch - */ - -template -std::shared_ptr> make_batch_k_query( - const raft::resources& res, - const cuvs::neighbors::brute_force::index& index, - raft::device_matrix_view query, - int64_t batch_size) -{ - return std::shared_ptr>( - new detail::gpu_batch_k_query(res, index, query, batch_size)); -} -} // namespace cuvs::neighbors::brute_force diff --git a/cpp/include/cuvs/neighbors/brute_force_types.hpp b/cpp/include/cuvs/neighbors/brute_force_types.hpp deleted file mode 100644 index 0d3252d71..000000000 --- a/cpp/include/cuvs/neighbors/brute_force_types.hpp +++ /dev/null @@ -1,283 +0,0 @@ -/* - * Copyright (c) 2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include "ann_types.hpp" -#include - -#include -#include -#include -#include -#include -#include -#include - -#include - -namespace cuvs::neighbors::brute_force { -/** - * @addtogroup brute_force_knn - * @{ - */ - -/** - * @brief Brute Force index. - * - * The index stores the dataset and norms for the dataset in device memory. - * - * @tparam T data element type - */ -template -struct index : ann::index { - public: - /** Distance metric used for retrieval */ - [[nodiscard]] constexpr inline cuvs::distance::DistanceType metric() const noexcept - { - return metric_; - } - - /** Total length of the index (number of vectors). */ - [[nodiscard]] constexpr inline int64_t size() const noexcept { return dataset_view_.extent(0); } - - /** Dimensionality of the data. */ - [[nodiscard]] constexpr inline uint32_t dim() const noexcept { return dataset_view_.extent(1); } - - /** Dataset [size, dim] */ - [[nodiscard]] inline auto dataset() const noexcept - -> raft::device_matrix_view - { - return dataset_view_; - } - - /** Dataset norms */ - [[nodiscard]] inline auto norms() const - -> raft::device_vector_view - { - return norms_view_.value(); - } - - /** Whether or not this index has dataset norms */ - [[nodiscard]] inline bool has_norms() const noexcept { return norms_view_.has_value(); } - - [[nodiscard]] inline T metric_arg() const noexcept { return metric_arg_; } - - // Don't allow copying the index for performance reasons (try avoiding copying data) - index(const index&) = delete; - index(index&&) = default; - auto operator=(const index&) -> index& = delete; - auto operator=(index&&) -> index& = default; - ~index() = default; - - /** Construct a brute force index from dataset - * - * Constructs a brute force index from a dataset. This lets us precompute norms for - * the dataset, providing a speed benefit over doing this at query time. - - * If the dataset is already in GPU memory, then this class stores a non-owning reference to - * the dataset. If the dataset is in host memory, it will be copied to the device and the - * index will own the device memory. - */ - template - index(raft::resources const& res, - raft::mdspan, raft::row_major, data_accessor> dataset, - std::optional>&& norms, - cuvs::distance::DistanceType metric, - T metric_arg = 0.0) - : ann::index(), - metric_(metric), - dataset_(raft::make_device_matrix(res, 0, 0)), - norms_(std::move(norms)), - metric_arg_(metric_arg) - { - if (norms_) { norms_view_ = raft::make_const_mdspan(norms_.value().view()); } - update_dataset(res, dataset); - raft::resource::sync_stream(res); - } - - /** Construct a brute force index from dataset - * - * This class stores a non-owning reference to the dataset and norms here. - * Having precomputed norms gives us a performance advantage at query time. - */ - index(raft::resources const& res, - raft::device_matrix_view dataset_view, - std::optional> norms_view, - cuvs::distance::DistanceType metric, - T metric_arg = 0.0) - : ann::index(), - metric_(metric), - dataset_(raft::make_device_matrix(res, 0, 0)), - dataset_view_(dataset_view), - norms_view_(norms_view), - metric_arg_(metric_arg) - { - } - - private: - /** - * Replace the dataset with a new dataset. - */ - void update_dataset(raft::resources const& res, - raft::device_matrix_view dataset) - { - dataset_view_ = dataset; - } - - /** - * Replace the dataset with a new dataset. - * - * We create a copy of the dataset on the device. The index manages the lifetime of this copy. - */ - void update_dataset(raft::resources const& res, - raft::host_matrix_view dataset) - { - dataset_ = raft::make_device_matrix(dataset.extents(0), dataset.extents(1)); - raft::copy(dataset_.data_handle(), - dataset.data_handle(), - dataset.size(), - resource::get_cuda_stream(res)); - dataset_view_ = raft::make_const_mdspan(dataset_.view()); - } - - cuvs::distance::DistanceType metric_; - raft::device_matrix dataset_; - std::optional> norms_; - std::optional> norms_view_; - raft::device_matrix_view dataset_view_; - T metric_arg_; -}; - -/** - * @brief Interface for performing queries over values of k - * - * This interface lets you iterate over batches of k from a brute_force::index. - * This lets you do things like retrieve the first 100 neighbors for a query, - * apply post processing to remove any unwanted items and then if needed get the - * next 100 closest neighbors for the query. - * - * This query interface exposes C++ iterators through the ::begin and ::end, and - * is compatible with range based for loops. - * - * Note that this class is an abstract class without any cuda dependencies, meaning - * that it doesn't require a cuda compiler to use - but also means it can't be directly - * instantiated. See the cuvs::neighbors::brute_force::make_batch_k_query - * function for usage examples. - * - * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset - */ -template -class batch_k_query { - public: - batch_k_query(const raft::resources& res, - int64_t index_size, - int64_t query_size, - int64_t batch_size) - : res(res), index_size(index_size), query_size(query_size), batch_size(batch_size) - { - } - virtual ~batch_k_query() {} - - using value_type = cuvs::neighbors::batch; - - class iterator { - public: - using value_type = cuvs::neighbors::batch; - using reference = const value_type&; - using pointer = const value_type*; - - iterator(const batch_k_query* query, int64_t offset = 0) - : current(query->res, 0, 0), batches(query->res, 0, 0), query(query), offset(offset) - { - query->load_batch(offset, query->batch_size, &batches); - query->slice_batch(batches, offset, query->batch_size, ¤t); - } - - reference operator*() const { return current; } - - pointer operator->() const { return ¤t; } - - iterator& operator++() - { - advance(query->batch_size); - return *this; - } - - iterator operator++(int) - { - iterator previous(*this); - operator++(); - return previous; - } - - /** - * @brief Advance the iterator, using a custom size for the next batch - * - * Using operator++ means that we will load up the same batch_size for each - * batch. This method allows us to get around this restriction, and load up - * arbitrary batch sizes on each iteration. - * See cuvs::neighbors::brute_force::make_batch_k_query for a usage example. - * - * @param[in] next_batch_size: size of the next batch to load up - */ - void advance(int64_t next_batch_size) - { - offset = std::min(offset + current.batch_size(), query->index_size); - if (offset + next_batch_size > batches.batch_size()) { - query->load_batch(offset, next_batch_size, &batches); - } - query->slice_batch(batches, offset, next_batch_size, ¤t); - } - - friend bool operator==(const iterator& lhs, const iterator& rhs) - { - return (lhs.query == rhs.query) && (lhs.offset == rhs.offset); - }; - friend bool operator!=(const iterator& lhs, const iterator& rhs) { return !(lhs == rhs); }; - - protected: - // the current batch of data - value_type current; - - // the currently loaded group of data (containing multiple batches of data that we can iterate - // through) - value_type batches; - - const batch_k_query* query; - int64_t offset, current_batch_size; - }; - - iterator begin() const { return iterator(this); } - iterator end() const { return iterator(this, index_size); } - - protected: - // these two methods need cuda code, and are implemented in the subclass - virtual void load_batch(int64_t offset, - int64_t next_batch_size, - batch* output) const = 0; - virtual void slice_batch(const value_type& input, - int64_t offset, - int64_t batch_size, - value_type* output) const = 0; - - const raft::resources& res; - int64_t index_size, query_size, batch_size; -}; -/** @} */ - -} // namespace cuvs::neighbors::brute_force diff --git a/cpp/include/cuvs/neighbors/cagra.hpp b/cpp/include/cuvs/neighbors/cagra.hpp index 8a4a8f017..3a0c60b78 100644 --- a/cpp/include/cuvs/neighbors/cagra.hpp +++ b/cpp/include/cuvs/neighbors/cagra.hpp @@ -18,11 +18,11 @@ #include "ann_types.hpp" #include -#include #include #include #include #include +#include #include namespace cuvs::neighbors::cagra { @@ -53,16 +53,17 @@ struct index_params : ann::index_params { size_t nn_descent_niter = 20; /** Build a raft CAGRA index params from an existing cuvs CAGRA index params. */ - operator raft::neighbors::cagra::index_params() const { + operator raft::neighbors::cagra::index_params() const + { return raft::neighbors::cagra::index_params{ { - .metric = static_cast((int)this->metric), - .metric_arg = this->metric_arg, + .metric = static_cast((int)this->metric), + .metric_arg = this->metric_arg, .add_data_on_build = this->add_data_on_build, }, .intermediate_graph_degree = intermediate_graph_degree, - .graph_degree = graph_degree, - .build_algo = static_cast((int)build_algo), + .graph_degree = graph_degree, + .build_algo = static_cast((int)build_algo), .nn_descent_niter = nn_descent_niter}; } }; @@ -122,7 +123,8 @@ struct search_params : ann::search_params { uint64_t rand_xor_mask = 0x128394; /** Build a raft CAGRA search params from an existing cuvs CAGRA search params. */ - operator raft::neighbors::cagra::search_params() const { + operator raft::neighbors::cagra::search_params() const + { raft::neighbors::cagra::search_params result = { {}, max_queries, @@ -156,7 +158,6 @@ static_assert(std::is_aggregate_v); */ template struct index : ann::index { - /** Build a cuvs CAGRA index from an existing RAFT CAGRA index. */ index(raft::neighbors::cagra::index&& raft_idx) : ann::index(), @@ -174,10 +175,7 @@ struct index : ann::index { } /** Total length of the index (number of vectors). */ - [[nodiscard]] constexpr inline auto size() const noexcept -> IdxT - { - return raft_index_->size(); - } + [[nodiscard]] constexpr inline auto size() const noexcept -> IdxT { return raft_index_->size(); } /** Dimensionality of the data. */ [[nodiscard]] constexpr inline auto dim() const noexcept -> uint32_t @@ -215,7 +213,8 @@ struct index : ann::index { index(raft::resources const& res, cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2Expanded) : ann::index(), - raft_index_(std::make_unique>(res, static_cast((int)metric))) + raft_index_(std::make_unique>( + res, static_cast((int)metric))) { } /** Construct an index from dataset and knn_graph arrays @@ -339,57 +338,55 @@ struct index : ann::index { { return raft_index_.get(); } - auto get_raft_index() -> raft::neighbors::cagra::index* - { - return raft_index_.get(); - } + auto get_raft_index() -> raft::neighbors::cagra::index* { return raft_index_.get(); } + private: std::unique_ptr> raft_index_; }; // Using device and host_matrix_view avoids needing to typedef multiple mdspans based on accessors -#define CUVS_INST_CAGRA_FUNCS(T, IdxT) \ - auto build(raft::resources const& handle, \ - const cuvs::neighbors::cagra::index_params& params, \ +#define CUVS_INST_CAGRA_FUNCS(T, IdxT) \ + auto build(raft::resources const& handle, \ + const cuvs::neighbors::cagra::index_params& params, \ raft::device_matrix_view dataset) \ - ->cuvs::neighbors::cagra::index; \ - \ - auto build(raft::resources const& handle, \ - const cuvs::neighbors::cagra::index_params& params, \ + ->cuvs::neighbors::cagra::index; \ + \ + auto build(raft::resources const& handle, \ + const cuvs::neighbors::cagra::index_params& params, \ raft::host_matrix_view dataset) \ - ->cuvs::neighbors::cagra::index; \ - \ - void build_device(raft::resources const& handle, \ - const cuvs::neighbors::cagra::index_params& params, \ + ->cuvs::neighbors::cagra::index; \ + \ + void build_device(raft::resources const& handle, \ + const cuvs::neighbors::cagra::index_params& params, \ raft::device_matrix_view dataset, \ - cuvs::neighbors::cagra::index& idx); \ - \ - void build_host(raft::resources const& handle, \ - const cuvs::neighbors::cagra::index_params& params, \ + cuvs::neighbors::cagra::index& idx); \ + \ + void build_host(raft::resources const& handle, \ + const cuvs::neighbors::cagra::index_params& params, \ raft::host_matrix_view dataset, \ - cuvs::neighbors::cagra::index& idx); \ - \ - void search(raft::resources const& handle, \ - cuvs::neighbors::cagra::search_params const& params, \ - const cuvs::neighbors::cagra::index& index, \ + cuvs::neighbors::cagra::index& idx); \ + \ + void search(raft::resources const& handle, \ + cuvs::neighbors::cagra::search_params const& params, \ + const cuvs::neighbors::cagra::index& index, \ raft::device_matrix_view queries, \ raft::device_matrix_view neighbors, \ raft::device_matrix_view distances); \ - void serialize_file(raft::resources const& handle, \ - const std::string& filename, \ - const cuvs::neighbors::cagra::index& index, \ - bool include_dataset = true); \ - \ - void deserialize_file(raft::resources const& handle, \ - const std::string& filename, \ - cuvs::neighbors::cagra::index* index); \ - void serialize(raft::resources const& handle, \ - std::string& str, \ - const cuvs::neighbors::cagra::index& index, \ - bool include_dataset = true); \ - \ - void deserialize(raft::resources const& handle, \ - const std::string& str, \ + void serialize_file(raft::resources const& handle, \ + const std::string& filename, \ + const cuvs::neighbors::cagra::index& index, \ + bool include_dataset = true); \ + \ + void deserialize_file(raft::resources const& handle, \ + const std::string& filename, \ + cuvs::neighbors::cagra::index* index); \ + void serialize(raft::resources const& handle, \ + std::string& str, \ + const cuvs::neighbors::cagra::index& index, \ + bool include_dataset = true); \ + \ + void deserialize(raft::resources const& handle, \ + const std::string& str, \ cuvs::neighbors::cagra::index* index); CUVS_INST_CAGRA_FUNCS(float, uint32_t); @@ -398,12 +395,12 @@ CUVS_INST_CAGRA_FUNCS(uint8_t, uint32_t); #undef CUVS_INST_CAGRA_FUNCS -#define CUVS_INST_CAGRA_OPTIMIZE(IdxT) \ - void optimize_device(raft::resources const& res, \ +#define CUVS_INST_CAGRA_OPTIMIZE(IdxT) \ + void optimize_device(raft::resources const& res, \ raft::device_matrix_view knn_graph, \ raft::host_matrix_view new_graph); \ - \ - void optimize_host(raft::resources const& res, \ + \ + void optimize_host(raft::resources const& res, \ raft::host_matrix_view knn_graph, \ raft::host_matrix_view new_graph); diff --git a/cpp/include/cuvs/neighbors/cagra_types.hpp b/cpp/include/cuvs/neighbors/cagra_types.hpp deleted file mode 100644 index 546279de1..000000000 --- a/cpp/include/cuvs/neighbors/cagra_types.hpp +++ /dev/null @@ -1,363 +0,0 @@ -/* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include "ann_types.hpp" -#include -#include - -#include -//#include -#include -#include -#include -#include -#include -#include - -#include -#include -#include -#include -#include - -#include -namespace cuvs::neighbors::cagra { -/** - * @addtogroup cagra - * @{ - */ - -/** - * @brief ANN algorithm used by CAGRA to build knn graph - * - */ -enum class graph_build_algo { - /* Use IVF-PQ to build all-neighbors knn graph */ - IVF_PQ, - /* Experimental, use NN-Descent to build all-neighbors knn graph */ - NN_DESCENT -}; - -struct index_params : ann::index_params { - /** Degree of input graph for pruning. */ - size_t intermediate_graph_degree = 128; - /** Degree of output graph. */ - size_t graph_degree = 64; - /** ANN algorithm to build knn graph. */ - graph_build_algo build_algo = graph_build_algo::IVF_PQ; - /** Number of Iterations to run if building with NN_DESCENT */ - size_t nn_descent_niter = 20; - - /** Build a raft CAGRA index params from an existing cuvs CAGRA index params. */ - operator raft::neighbors::cagra::index_params() const { - return raft::neighbors::cagra::index_params{ - { - .metric = static_cast((int)this->metric), - .metric_arg = this->metric_arg, - .add_data_on_build = this->add_data_on_build, - }, - .intermediate_graph_degree = intermediate_graph_degree, - .graph_degree = graph_degree, - .build_algo = static_cast((int)build_algo), - .nn_descent_niter = nn_descent_niter}; - } -}; - -enum class search_algo { - /** For large batch sizes. */ - SINGLE_CTA, - /** For small batch sizes. */ - MULTI_CTA, - MULTI_KERNEL, - AUTO -}; - -enum class hash_mode { HASH, SMALL, AUTO }; - -struct search_params : ann::search_params { - /** Maximum number of queries to search at the same time (batch size). Auto select when 0.*/ - size_t max_queries = 0; - - /** Number of intermediate search results retained during the search. - * - * This is the main knob to adjust trade off between accuracy and search speed. - * Higher values improve the search accuracy. - */ - size_t itopk_size = 64; - - /** Upper limit of search iterations. Auto select when 0.*/ - size_t max_iterations = 0; - - // In the following we list additional search parameters for fine tuning. - // Reasonable default values are automatically chosen. - - /** Which search implementation to use. */ - search_algo algo = search_algo::AUTO; - - /** Number of threads used to calculate a single distance. 4, 8, 16, or 32. */ - size_t team_size = 0; - - /** Number of graph nodes to select as the starting point for the search in each iteration. aka - * search width?*/ - size_t search_width = 1; - /** Lower limit of search iterations. */ - size_t min_iterations = 0; - - /** Thread block size. 0, 64, 128, 256, 512, 1024. Auto selection when 0. */ - size_t thread_block_size = 0; - /** Hashmap type. Auto selection when AUTO. */ - hash_mode hashmap_mode = hash_mode::AUTO; - /** Lower limit of hashmap bit length. More than 8. */ - size_t hashmap_min_bitlen = 0; - /** Upper limit of hashmap fill rate. More than 0.1, less than 0.9.*/ - float hashmap_max_fill_rate = 0.5; - - /** Number of iterations of initial random seed node selection. 1 or more. */ - uint32_t num_random_samplings = 1; - /** Bit mask used for initial random seed node selection. */ - uint64_t rand_xor_mask = 0x128394; - - /** Build a raft CAGRA search params from an existing cuvs CAGRA search params. */ - operator raft::neighbors::cagra::search_params() const { - raft::neighbors::cagra::search_params result = { - {}, - max_queries, - itopk_size, - max_iterations, - static_cast((int)algo), - team_size, - search_width, - min_iterations, - thread_block_size, - static_cast((int)hashmap_mode), - hashmap_min_bitlen, - hashmap_max_fill_rate, - num_random_samplings, - rand_xor_mask}; - return result; - } -}; - -static_assert(std::is_aggregate_v); -static_assert(std::is_aggregate_v); - -/** - * @brief CAGRA index. - * - * The index stores the dataset and a kNN graph in device memory. - * - * @tparam T data element type - * @tparam IdxT type of the vector indices (represent dataset.extent(0)) - * - */ -template -struct index : ann::index { - - /** Build a cuvs CAGRA index from an existing RAFT CAGRA index. */ - index(raft::neighbors::cagra::index&& raft_idx) - : ann::index(), - raft_index_{std::make_unique>(std::move(raft_idx))} - { - } - static_assert(!raft::is_narrowing_v, - "IdxT must be able to represent all values of uint32_t"); - - public: - /** Distance metric used for clustering. */ - [[nodiscard]] constexpr inline auto metric() const noexcept -> cuvs::distance::DistanceType - { - return static_cast((int)raft_index_->metric()); - } - - /** Total length of the index (number of vectors). */ - [[nodiscard]] constexpr inline auto size() const noexcept -> IdxT - { - return raft_index_->size(); - } - - /** Dimensionality of the data. */ - [[nodiscard]] constexpr inline auto dim() const noexcept -> uint32_t - { - return raft_index_->dim(); - } - /** Graph degree */ - [[nodiscard]] constexpr inline auto graph_degree() const noexcept -> uint32_t - { - return raft_index_->graph_degree(); - } - - /** Dataset [size, dim] */ - [[nodiscard]] inline auto dataset() const noexcept - -> raft::device_matrix_view - { - return raft_index_->dataset(); - } - - /** neighborhood graph [size, graph-degree] */ - [[nodiscard]] inline auto graph() const noexcept - -> raft::device_matrix_view - { - return raft_index_->graph(); - } - - // Don't allow copying the index for performance reasons (try avoiding copying data) - index(const index&) = delete; - index(index&&) = default; - auto operator=(const index&) -> index& = delete; - auto operator=(index&&) -> index& = default; - ~index() = default; - - /** Construct an empty index. */ - index(raft::resources const& res, - cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2Expanded) - : ann::index(), - raft_index_(std::make_unique>(res, static_cast((int)metric))) - { - } - /** Construct an index from dataset and knn_graph arrays - * - * If the dataset and graph is already in GPU memory, then the index is just a thin wrapper around - * these that stores a non-owning a reference to the arrays. - * - * The constructor also accepts host arrays. In that case they are copied to the device, and the - * device arrays will be owned by the index. - * - * In case the dasates rows are not 16 bytes aligned, then we create a padded copy in device - * memory to ensure alignment for vectorized load. - * - * Usage examples: - * - * - Cagra index is normally created by the cagra::build - * @code{.cpp} - * using namespace cuvs::neighbors::experimental; - * auto dataset = raft::make_host_matrix(n_rows, n_cols); - * load_dataset(dataset.view()); - * // use default index parameters - * cagra::index_params index_params; - * // create and fill the index from a [N, D] dataset - * auto index = cagra::build(res, index_params, dataset); - * // use default search parameters - * cagra::search_params search_params; - * // search K nearest neighbours - * auto neighbors = raft::make_device_matrix(res, n_queries, k); - * auto distances = raft::make_device_matrix(res, n_queries, k); - * cagra::search(res, search_params, index, queries, neighbors, distances); - * @endcode - * In the above example, we have passed a host dataset to build. The returned index will own a - * device copy of the dataset and the knn_graph. In contrast, if we pass the dataset as a - * raft::device_mdspan to build, then it will only store a reference to it. - * - * - Constructing index using existing knn-graph - * @code{.cpp} - * using namespace cuvs::neighbors::experimental; - * - * auto dataset = raft::make_device_matrix(res, n_rows, n_cols); - * auto knn_graph = raft::make_device_matrix(res, n_rows, graph_degree); - * - * // custom loading and graph creation - * // load_dataset(dataset.view()); - * // create_knn_graph(knn_graph.view()); - * - * // Wrap the existing device arrays into an index structure - * cagra::index index(res, metric, raft::make_const_mdspan(dataset.view()), - * raft::make_const_mdspan(knn_graph.view())); - * - * // Both knn_graph and dataset objects have to be in scope while the index is used because - * // the index only stores a reference to these. - * cagra::search(res, search_params, index, queries, neighbors, distances); - * @endcode - * - */ - template - index(raft::resources const& res, - cuvs::distance::DistanceType metric, - raft::mdspan, raft::row_major, data_accessor> dataset, - raft::mdspan, raft::row_major, graph_accessor> - knn_graph) - : ann::index(), - raft_index_(std::make_unique>( - res, static_cast((int)metric), dataset, knn_graph)) - { - RAFT_EXPECTS(dataset.extent(0) == knn_graph.extent(0), - "Dataset and knn_graph must have equal number of rows"); - update_dataset(res, dataset); - update_graph(res, knn_graph); - raft::resource::sync_stream(res); - } - - /** - * Replace the dataset with a new dataset. - * - * If the new dataset rows are aligned on 16 bytes, then only a reference is stored to the - * dataset. It is the caller's responsibility to ensure that dataset stays alive as long as the - * index. - */ - void update_dataset(raft::resources const& res, - raft::device_matrix_view dataset) - { - raft_index_->update_dataset(res, dataset); - } - /** - * Replace the dataset with a new dataset. - * - * We create a copy of the dataset on the device. The index manages the lifetime of this copy. - */ - void update_dataset(raft::resources const& res, - raft::host_matrix_view dataset) - { - raft_index_->update_dataset(res, dataset); - } - - /** - * Replace the graph with a new graph. - * - * Since the new graph is a device array, we store a reference to that, and it is - * the caller's responsibility to ensure that knn_graph stays alive as long as the index. - */ - void update_graph(raft::resources const& res, - raft::device_matrix_view knn_graph) - { - raft_index_->update_graph(res, knn_graph); - } - - /** - * Replace the graph with a new graph. - * - * We create a copy of the graph on the device. The index manages the lifetime of this copy. - */ - void update_graph(raft::resources const& res, - raft::host_matrix_view knn_graph) - { - raft_index_->update_graph(res, knn_graph); - } - - auto get_raft_index() const -> const raft::neighbors::cagra::index* - { - return raft_index_.get(); - } - auto get_raft_index() -> raft::neighbors::cagra::index* - { - return raft_index_.get(); - } - private: - std::unique_ptr> raft_index_; -}; - -/** @} */ - -} // namespace cuvs::neighbors::cagra diff --git a/cpp/include/cuvs/neighbors/epsilon_neighborhood.cuh b/cpp/include/cuvs/neighbors/epsilon_neighborhood.cuh deleted file mode 100644 index dfa300c22..000000000 --- a/cpp/include/cuvs/neighbors/epsilon_neighborhood.cuh +++ /dev/null @@ -1,123 +0,0 @@ -/* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#ifndef __EPSILON_NEIGH_H -#define __EPSILON_NEIGH_H - -#pragma once - -#include -#include -#include -#include - -namespace cuvs::neighbors::epsilon_neighborhood { - -/** - * @brief Computes epsilon neighborhood for the L2-Squared distance metric - * - * @tparam value_t IO and math type - * @tparam idx_t Index type - * - * @param[out] adj adjacency matrix [row-major] [on device] [dim = m x n] - * @param[out] vd vertex degree array [on device] [len = m + 1] - * `vd + m` stores the total number of edges in the adjacency - * matrix. Pass a nullptr if you don't need this info. - * @param[in] x first matrix [row-major] [on device] [dim = m x k] - * @param[in] y second matrix [row-major] [on device] [dim = n x k] - * @param[in] m number of rows in x - * @param[in] n number of rows in y - * @param[in] k number of columns in x and k - * @param[in] eps defines epsilon neighborhood radius (should be passed as - * squared as we compute L2-squared distance in this method) - * @param[in] stream cuda stream - */ -template -void epsUnexpL2SqNeighborhood(bool* adj, - idx_t* vd, - const value_t* x, - const value_t* y, - idx_t m, - idx_t n, - idx_t k, - value_t eps, - cudaStream_t stream) -{ - spatial::knn::detail::epsUnexpL2SqNeighborhood( - adj, vd, x, y, m, n, k, eps, stream); -} - -/** - * @defgroup epsilon_neighbors Epislon Neighborhood Operations - * @{ - */ - -/** - * @brief Computes epsilon neighborhood for the L2-Squared distance metric and given ball size. - * The epsilon neighbors is represented by a dense boolean adjacency matrix of size m * n and - * an array of degrees for each vertex, which can be used as a compressed sparse row (CSR) - * indptr array. - * - * @code{.cpp} - * #include - * #include - * #include - * using namespace cuvs::neighbors; - * raft::raft::resources handle; - * ... - * auto adj = raft::make_device_matrix(handle, m * n); - * auto vd = raft::make_device_vector(handle, m+1); - * epsilon_neighborhood::eps_neighbors_l2sq(handle, x, y, adj.view(), vd.view(), eps); - * @endcode - * - * @tparam value_t IO and math type - * @tparam idx_t Index type - * @tparam matrix_idx_t matrix indexing type - * - * @param[in] handle raft handle to manage library resources - * @param[in] x first matrix [row-major] [on device] [dim = m x k] - * @param[in] y second matrix [row-major] [on device] [dim = n x k] - * @param[out] adj adjacency matrix [row-major] [on device] [dim = m x n] - * @param[out] vd vertex degree array [on device] [len = m + 1] - * `vd + m` stores the total number of edges in the adjacency - * matrix. Pass a nullptr if you don't need this info. - * @param[in] eps defines epsilon neighborhood radius (should be passed as - * squared as we compute L2-squared distance in this method) - */ -template -void eps_neighbors_l2sq(raft::resources const& handle, - raft::device_matrix_view x, - raft::device_matrix_view y, - raft::device_matrix_view adj, - raft::device_vector_view vd, - value_t eps) -{ - epsUnexpL2SqNeighborhood(adj.data_handle(), - vd.data_handle(), - x.data_handle(), - y.data_handle(), - x.extent(0), - y.extent(0), - x.extent(1), - eps, - resource::get_cuda_stream(handle)); -} - -/** @} */ // end group epsilon_neighbors - -} // namespace cuvs::neighbors::epsilon_neighborhood - -#endif \ No newline at end of file diff --git a/cpp/include/cuvs/neighbors/ivf_flat-ext.cuh b/cpp/include/cuvs/neighbors/ivf_flat-ext.cuh deleted file mode 100644 index 3b66a589b..000000000 --- a/cpp/include/cuvs/neighbors/ivf_flat-ext.cuh +++ /dev/null @@ -1,206 +0,0 @@ -/* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include // int64_t - -#include -#include // cuvs::neighbors::ivf_flat::index -#include // raft::device_matrix_view -#include // raft::resources -#include // RAFT_EXPLICIT -#include // rmm::mr::device_memory_resource - -#ifdef RAFT_EXPLICIT_INSTANTIATE_ONLY - -namespace cuvs::neighbors::ivf_flat { - -template -auto build(raft::resources const& handle, - const index_params& params, - const T* dataset, - IdxT n_rows, - uint32_t dim) -> index RAFT_EXPLICIT; - -template -auto build(raft::resources const& handle, - const index_params& params, - raft::device_matrix_view dataset) - -> index RAFT_EXPLICIT; - -template -void build(raft::resources const& handle, - const index_params& params, - raft::device_matrix_view dataset, - cuvs::neighbors::ivf_flat::index& idx) RAFT_EXPLICIT; - -template -auto extend(raft::resources const& handle, - const index& orig_index, - const T* new_vectors, - const IdxT* new_indices, - IdxT n_rows) -> index RAFT_EXPLICIT; - -template -auto extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, - std::optional> new_indices, - const index& orig_index) -> index RAFT_EXPLICIT; - -template -void extend(raft::resources const& handle, - index* index, - const T* new_vectors, - const IdxT* new_indices, - IdxT n_rows) RAFT_EXPLICIT; - -template -void extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, - std::optional> new_indices, - index* index) RAFT_EXPLICIT; - -template -void search_with_filtering(raft::resources const& handle, - const search_params& params, - const index& index, - const T* queries, - uint32_t n_queries, - uint32_t k, - IdxT* neighbors, - float* distances, - rmm::mr::device_memory_resource* mr = nullptr, - IvfSampleFilterT sample_filter = IvfSampleFilterT()) RAFT_EXPLICIT; - -template -void search(raft::resources const& handle, - const search_params& params, - const index& index, - const T* queries, - uint32_t n_queries, - uint32_t k, - IdxT* neighbors, - float* distances, - rmm::mr::device_memory_resource* mr = nullptr) RAFT_EXPLICIT; - -template -void search_with_filtering(raft::resources const& handle, - const search_params& params, - const index& index, - raft::device_matrix_view queries, - raft::device_matrix_view neighbors, - raft::device_matrix_view distances, - IvfSampleFilterT sample_filter = IvfSampleFilterT()) RAFT_EXPLICIT; - -template -void search(raft::resources const& handle, - const search_params& params, - const index& index, - raft::device_matrix_view queries, - raft::device_matrix_view neighbors, - raft::device_matrix_view distances) RAFT_EXPLICIT; - -} // namespace cuvs::neighbors::ivf_flat - -#endif // RAFT_EXPLICIT_INSTANTIATE_ONLY - -#define instantiate_raft_neighbors_ivf_flat_build(T, IdxT) \ - extern template auto cuvs::neighbors::ivf_flat::build( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_flat::index_params& params, \ - const T* dataset, \ - IdxT n_rows, \ - uint32_t dim) \ - ->cuvs::neighbors::ivf_flat::index; \ - \ - extern template auto cuvs::neighbors::ivf_flat::build( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_flat::index_params& params, \ - raft::device_matrix_view dataset) \ - ->cuvs::neighbors::ivf_flat::index; \ - \ - extern template void cuvs::neighbors::ivf_flat::build( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_flat::index_params& params, \ - raft::device_matrix_view dataset, \ - cuvs::neighbors::ivf_flat::index& idx); - -instantiate_raft_neighbors_ivf_flat_build(float, int64_t); -instantiate_raft_neighbors_ivf_flat_build(int8_t, int64_t); -instantiate_raft_neighbors_ivf_flat_build(uint8_t, int64_t); -#undef instantiate_raft_neighbors_ivf_flat_build - -#define instantiate_raft_neighbors_ivf_flat_extend(T, IdxT) \ - extern template auto cuvs::neighbors::ivf_flat::extend( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_flat::index& orig_index, \ - const T* new_vectors, \ - const IdxT* new_indices, \ - IdxT n_rows) \ - ->cuvs::neighbors::ivf_flat::index; \ - \ - extern template auto cuvs::neighbors::ivf_flat::extend( \ - raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ - std::optional> new_indices, \ - const cuvs::neighbors::ivf_flat::index& orig_index) \ - ->cuvs::neighbors::ivf_flat::index; \ - \ - extern template void cuvs::neighbors::ivf_flat::extend( \ - raft::resources const& handle, \ - cuvs::neighbors::ivf_flat::index* index, \ - const T* new_vectors, \ - const IdxT* new_indices, \ - IdxT n_rows); \ - \ - extern template void cuvs::neighbors::ivf_flat::extend( \ - raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ - std::optional> new_indices, \ - cuvs::neighbors::ivf_flat::index* index); - -instantiate_raft_neighbors_ivf_flat_extend(float, int64_t); -instantiate_raft_neighbors_ivf_flat_extend(int8_t, int64_t); -instantiate_raft_neighbors_ivf_flat_extend(uint8_t, int64_t); - -#undef instantiate_raft_neighbors_ivf_flat_extend - -#define instantiate_raft_neighbors_ivf_flat_search(T, IdxT) \ - extern template void cuvs::neighbors::ivf_flat::search( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_flat::search_params& params, \ - const cuvs::neighbors::ivf_flat::index& index, \ - const T* queries, \ - uint32_t n_queries, \ - uint32_t k, \ - IdxT* neighbors, \ - float* distances, \ - rmm::mr::device_memory_resource* mr); \ - \ - extern template void cuvs::neighbors::ivf_flat::search( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_flat::search_params& params, \ - const cuvs::neighbors::ivf_flat::index& index, \ - raft::device_matrix_view queries, \ - raft::device_matrix_view neighbors, \ - raft::device_matrix_view distances); - -instantiate_raft_neighbors_ivf_flat_search(float, int64_t); -instantiate_raft_neighbors_ivf_flat_search(int8_t, int64_t); -instantiate_raft_neighbors_ivf_flat_search(uint8_t, int64_t); - -#undef instantiate_raft_neighbors_ivf_flat_search diff --git a/cpp/include/cuvs/neighbors/ivf_flat-inl.cuh b/cpp/include/cuvs/neighbors/ivf_flat-inl.cuh deleted file mode 100644 index d956f060c..000000000 --- a/cpp/include/cuvs/neighbors/ivf_flat-inl.cuh +++ /dev/null @@ -1,602 +0,0 @@ -/* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include -#include -#include -#include - -#include - -#include -#include -#include - -namespace cuvs::neighbors::ivf_flat { - -/** - * @brief Build the index from the dataset for efficient search. - * - * NB: Currently, the following distance metrics are supported: - * - L2Expanded - * - L2Unexpanded - * - InnerProduct - * - * Usage example: - * @code{.cpp} - * using namespace cuvs::neighbors; - * // use default index parameters - * ivf_flat::index_params index_params; - * // create and fill the index from a [N, D] dataset - * auto index = ivf_flat::build(handle, index_params, dataset, N, D); - * // use default search parameters - * ivf_flat::search_params search_params; - * // search K nearest neighbours for each of the N queries - * ivf_flat::search(handle, search_params, index, queries, N, K, out_inds, out_dists); - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset - * - * @param[in] handle - * @param[in] params configure the index building - * @param[in] dataset a device pointer to a row-major matrix [n_rows, dim] - * @param[in] n_rows the number of samples - * @param[in] dim the dimensionality of the data - * - * @return the constructed ivf-flat index - */ -template -auto build(raft::resources const& handle, - const index_params& params, - const T* dataset, - IdxT n_rows, - uint32_t dim) -> index -{ - return cuvs::neighbors::ivf_flat::detail::build(handle, params, dataset, n_rows, dim); -} - -/** - * @defgroup ivf_flat IVF Flat Algorithm - * @{ - */ - -/** - * @brief Build the index from the dataset for efficient search. - * - * NB: Currently, the following distance metrics are supported: - * - L2Expanded - * - L2Unexpanded - * - InnerProduct - * - * Usage example: - * @code{.cpp} - * using namespace cuvs::neighbors; - * // use default index parameters - * ivf_flat::index_params index_params; - * // create and fill the index from a [N, D] dataset - * auto index = ivf_flat::build(handle, dataset, index_params); - * // use default search parameters - * ivf_flat::search_params search_params; - * // search K nearest neighbours for each of the N queries - * ivf_flat::search(handle, search_params, index, queries, out_inds, out_dists); - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset - * - * @param[in] handle - * @param[in] params configure the index building - * @param[in] dataset a device pointer to a row-major matrix [n_rows, dim] - * - * @return the constructed ivf-flat index - */ -template -auto build(raft::resources const& handle, - const index_params& params, - raft::device_matrix_view dataset) -> index -{ - return cuvs::neighbors::ivf_flat::detail::build(handle, - params, - dataset.data_handle(), - static_cast(dataset.extent(0)), - static_cast(dataset.extent(1))); -} - -/** - * @brief Build the index from the dataset for efficient search. - * - * NB: Currently, the following distance metrics are supported: - * - L2Expanded - * - L2Unexpanded - * - InnerProduct - * - * Usage example: - * @code{.cpp} - * using namespace cuvs::neighbors; - * // use default index parameters - * ivf_flat::index_params index_params; - * // create and fill the index from a [N, D] dataset - * ivf_flat::index index; - * ivf_flat::build(handle, dataset, index_params, index); - * // use default search parameters - * ivf_flat::search_params search_params; - * // search K nearest neighbours for each of the N queries - * ivf_flat::search(handle, search_params, index, queries, out_inds, out_dists); - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset - * - * @param[in] handle - * @param[in] params configure the index building - * @param[in] dataset raft::device_matrix_view to a row-major matrix [n_rows, dim] - * @param[out] idx reference to ivf_flat::index - * - */ -template -void build(raft::resources const& handle, - const index_params& params, - raft::device_matrix_view dataset, - cuvs::neighbors::ivf_flat::index& idx) -{ - idx = cuvs::neighbors::ivf_flat::detail::build(handle, - params, - dataset.data_handle(), - static_cast(dataset.extent(0)), - static_cast(dataset.extent(1))); -} - -/** @} */ - -/** - * @brief Build a new index containing the data of the original plus new extra vectors. - * - * Implementation note: - * The new data is clustered according to existing kmeans clusters, then the cluster - * centers are adjusted to match the newly labeled data. - * - * Usage example: - * @code{.cpp} - * using namespace cuvs::neighbors; - * ivf_flat::index_params index_params; - * index_params.add_data_on_build = false; // don't populate index on build - * index_params.kmeans_trainset_fraction = 1.0; // use whole dataset for kmeans training - * // train the index from a [N, D] dataset - * auto index_empty = ivf_flat::build(handle, index_params, dataset, N, D); - * // fill the index with the data - * auto index = ivf_flat::extend(handle, index_empty, dataset, nullptr, N); - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset - * - * @param[in] handle - * @param[in] orig_index original index - * @param[in] new_vectors a device pointer to a row-major matrix [n_rows, index.dim()] - * @param[in] new_indices a device pointer to a vector of indices [n_rows]. - * If the original index is empty (`orig_index.size() == 0`), you can pass `nullptr` - * here to imply a continuous range `[0...n_rows)`. - * @param[in] n_rows number of rows in `new_vectors` - * - * @return the constructed extended ivf-flat index - */ -template -auto extend(raft::resources const& handle, - const index& orig_index, - const T* new_vectors, - const IdxT* new_indices, - IdxT n_rows) -> index -{ - return cuvs::neighbors::ivf_flat::detail::extend( - handle, orig_index, new_vectors, new_indices, n_rows); -} - -/** - * @ingroup ivf_flat - * @{ - */ - -/** - * @brief Build a new index containing the data of the original plus new extra vectors. - * - * Implementation note: - * The new data is clustered according to existing kmeans clusters, then the cluster - * centers are adjusted to match the newly labeled data. - * - * Usage example: - * @code{.cpp} - * using namespace cuvs::neighbors; - * ivf_flat::index_params index_params; - * index_params.add_data_on_build = false; // don't populate index on build - * index_params.kmeans_trainset_fraction = 1.0; // use whole dataset for kmeans training - * // train the index from a [N, D] dataset - * auto index_empty = ivf_flat::build(handle, dataset, index_params, dataset); - * // fill the index with the data - * std::optional> no_op = std::nullopt; - * auto index = ivf_flat::extend(handle, index_empty, no_op, dataset); - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset - * - * @param[in] handle - * @param[in] new_vectors raft::device_matrix_view to a row-major matrix [n_rows, index.dim()] - * @param[in] new_indices optional raft::device_vector_view to a vector of indices [n_rows]. - * If the original index is empty (`orig_index.size() == 0`), you can pass `std::nullopt` - * here to imply a continuous range `[0...n_rows)`. - * @param[in] orig_index original index - * - * @return the constructed extended ivf-flat index - */ -template -auto extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, - std::optional> new_indices, - const index& orig_index) -> index -{ - return extend(handle, - orig_index, - new_vectors.data_handle(), - new_indices.has_value() ? new_indices.value().data_handle() : nullptr, - new_vectors.extent(0)); -} - -/** @} */ - -/** - * @brief Extend the index in-place with the new data. - * - * Usage example: - * @code{.cpp} - * using namespace cuvs::neighbors; - * ivf_flat::index_params index_params; - * index_params.add_data_on_build = false; // don't populate index on build - * index_params.kmeans_trainset_fraction = 1.0; // use whole dataset for kmeans training - * // train the index from a [N, D] dataset - * auto index_empty = ivf_flat::build(handle, index_params, dataset, N, D); - * // fill the index with the data - * ivf_flat::extend(handle, index_empty, dataset, nullptr, N); - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset - * - * @param handle - * @param[inout] index - * @param[in] new_vectors a device pointer to a row-major matrix [n_rows, index.dim()] - * @param[in] new_indices a device pointer to a vector of indices [n_rows]. - * If the original index is empty (`orig_index.size() == 0`), you can pass `nullptr` - * here to imply a continuous range `[0...n_rows)`. - * @param[in] n_rows the number of samples - */ -template -void extend(raft::resources const& handle, - index* index, - const T* new_vectors, - const IdxT* new_indices, - IdxT n_rows) -{ - cuvs::neighbors::ivf_flat::detail::extend(handle, index, new_vectors, new_indices, n_rows); -} - -/** - * @ingroup ivf_flat - * @{ - */ - -/** - * @brief Extend the index in-place with the new data. - * - * Usage example: - * @code{.cpp} - * using namespace cuvs::neighbors; - * ivf_flat::index_params index_params; - * index_params.add_data_on_build = false; // don't populate index on build - * index_params.kmeans_trainset_fraction = 1.0; // use whole dataset for kmeans training - * // train the index from a [N, D] dataset - * auto index_empty = ivf_flat::build(handle, index_params, dataset); - * // fill the index with the data - * std::optional> no_op = std::nullopt; - * ivf_flat::extend(handle, dataset, no_opt, &index_empty); - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset - * - * @param[in] handle - * @param[in] new_vectors raft::device_matrix_view to a row-major matrix [n_rows, index.dim()] - * @param[in] new_indices optional raft::device_vector_view to a vector of indices [n_rows]. - * If the original index is empty (`orig_index.size() == 0`), you can pass `std::nullopt` - * here to imply a continuous range `[0...n_rows)`. - * @param[inout] index pointer to index, to be overwritten in-place - */ -template -void extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, - std::optional> new_indices, - index* index) -{ - extend(handle, - index, - new_vectors.data_handle(), - new_indices.has_value() ? new_indices.value().data_handle() : nullptr, - static_cast(new_vectors.extent(0))); -} - -/** @} */ - -/** - * @brief Search ANN using the constructed index with the given filter. - * - * See the [ivf_flat::build](#ivf_flat::build) documentation for a usage example. - * - * Note, this function requires a temporary buffer to store intermediate results between cuda kernel - * calls, which may lead to undesirable allocations and slowdown. To alleviate the problem, you can - * pass a pool memory resource or a large enough pre-allocated memory resource to reduce or - * eliminate entirely allocations happening within `search`: - * @code{.cpp} - * ... - * // Create a pooling memory resource with a pre-defined initial size. - * rmm::mr::pool_memory_resource mr( - * rmm::mr::get_current_device_resource(), 1024 * 1024); - * // use default search parameters - * ivf_flat::search_params search_params; - * filtering::none_ivf_sample_filter filter; - * // Use the same allocator across multiple searches to reduce the number of - * // cuda memory allocations - * ivf_flat::search_with_filtering( - * handle, search_params, index, queries1, N1, K, out_inds1, out_dists1, &mr, filter); - * ivf_flat::search_with_filtering( - * handle, search_params, index, queries2, N2, K, out_inds2, out_dists2, &mr, filter); - * ivf_flat::search_with_filtering( - * handle, search_params, index, queries3, N3, K, out_inds3, out_dists3, &mr, filter); - * ... - * @endcode - * The exact size of the temporary buffer depends on multiple factors and is an implementation - * detail. However, you can safely specify a small initial size for the memory pool, so that only a - * few allocations happen to grow it during the first invocations of the `search`. - * - * @tparam T data element type - * @tparam IdxT type of the indices - * @tparam IvfSampleFilterT Device filter function, with the signature - * `(uint32_t query_ix, uint32 cluster_ix, uint32_t sample_ix) -> bool` or - * `(uint32_t query_ix, uint32 sample_ix) -> bool` - * - * @param[in] handle - * @param[in] params configure the search - * @param[in] index ivf-flat constructed index - * @param[in] queries a device pointer to a row-major matrix [n_queries, index->dim()] - * @param[in] n_queries the batch size - * @param[in] k the number of neighbors to find for each query. - * @param[out] neighbors a device pointer to the indices of the neighbors in the source dataset - * [n_queries, k] - * @param[out] distances a device pointer to the distances to the selected neighbors [n_queries, k] - * @param[in] mr an optional memory resource to use across the searches (you can provide a large - * enough memory pool here to avoid memory allocations within search). - * @param[in] sample_filter a device filter function that greenlights samples for a given query - */ -template -void search_with_filtering(raft::resources const& handle, - const search_params& params, - const index& index, - const T* queries, - uint32_t n_queries, - uint32_t k, - IdxT* neighbors, - float* distances, - rmm::mr::device_memory_resource* mr = nullptr, - IvfSampleFilterT sample_filter = IvfSampleFilterT()) -{ - cuvs::neighbors::ivf_flat::detail::search( - handle, params, index, queries, n_queries, k, neighbors, distances, mr, sample_filter); -} - -/** - * @brief Search ANN using the constructed index using the given filter. - * - * See the [ivf_flat::build](#ivf_flat::build) documentation for a usage example. - * - * Note, this function requires a temporary buffer to store intermediate results between cuda kernel - * calls, which may lead to undesirable allocations and slowdown. To alleviate the problem, you can - * pass a pool memory resource or a large enough pre-allocated memory resource to reduce or - * eliminate entirely allocations happening within `search`: - * @code{.cpp} - * ... - * // Create a pooling memory resource with a pre-defined initial size. - * rmm::mr::pool_memory_resource mr( - * rmm::mr::get_current_device_resource(), 1024 * 1024); - * // use default search parameters - * ivf_flat::search_params search_params; - * // Use the same allocator across multiple searches to reduce the number of - * // cuda memory allocations - * ivf_flat::search(handle, search_params, index, queries1, N1, K, out_inds1, out_dists1, &mr); - * ivf_flat::search(handle, search_params, index, queries2, N2, K, out_inds2, out_dists2, &mr); - * ivf_flat::search(handle, search_params, index, queries3, N3, K, out_inds3, out_dists3, &mr); - * ... - * @endcode - * The exact size of the temporary buffer depends on multiple factors and is an implementation - * detail. However, you can safely specify a small initial size for the memory pool, so that only a - * few allocations happen to grow it during the first invocations of the `search`. - * - * @tparam T data element type - * @tparam IdxT type of the indices - * - * @param[in] handle - * @param[in] params configure the search - * @param[in] index ivf-flat constructed index - * @param[in] queries a device pointer to a row-major matrix [n_queries, index->dim()] - * @param[in] n_queries the batch size - * @param[in] k the number of neighbors to find for each query. - * @param[out] neighbors a device pointer to the indices of the neighbors in the source dataset - * [n_queries, k] - * @param[out] distances a device pointer to the distances to the selected neighbors [n_queries, k] - * @param[in] mr an optional memory resource to use across the searches (you can provide a large - * enough memory pool here to avoid memory allocations within search). - */ -template -void search(raft::resources const& handle, - const search_params& params, - const index& index, - const T* queries, - uint32_t n_queries, - uint32_t k, - IdxT* neighbors, - float* distances, - rmm::mr::device_memory_resource* mr = nullptr) -{ - cuvs::neighbors::ivf_flat::detail::search(handle, - params, - index, - queries, - n_queries, - k, - neighbors, - distances, - mr, - cuvs::neighbors::filtering::none_ivf_sample_filter()); -} - -/** - * @ingroup ivf_flat - * @{ - */ - -/** - * @brief Search ANN using the constructed index with the given filter. - * - * See the [ivf_flat::build](#ivf_flat::build) documentation for a usage example. - * - * Note, this function requires a temporary buffer to store intermediate results between cuda kernel - * calls, which may lead to undesirable allocations and slowdown. To alleviate the problem, you can - * pass a pool memory resource or a large enough pre-allocated memory resource to reduce or - * eliminate entirely allocations happening within `search`: - * @code{.cpp} - * ... - * // use default search parameters - * ivf_flat::search_params search_params; - * filtering::none_ivf_sample_filter filter; - * // Use the same allocator across multiple searches to reduce the number of - * // cuda memory allocations - * ivf_flat::search_with_filtering( - * handle, search_params, index, queries1, out_inds1, out_dists1, filter); - * ivf_flat::search_with_filtering( - * handle, search_params, index, queries2, out_inds2, out_dists2, filter); - * ivf_flat::search_with_filtering( - * handle, search_params, index, queries3, out_inds3, out_dists3, filter); - * ... - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices - * @tparam IvfSampleFilterT Device filter function, with the signature - * `(uint32_t query_ix, uint32 cluster_ix, uint32_t sample_ix) -> bool` or - * `(uint32_t query_ix, uint32 sample_ix) -> bool` - * - * @param[in] handle - * @param[in] params configure the search - * @param[in] index ivf-flat constructed index - * @param[in] queries a device pointer to a row-major matrix [n_queries, index->dim()] - * @param[out] neighbors a device pointer to the indices of the neighbors in the source dataset - * [n_queries, k] - * @param[out] distances a device pointer to the distances to the selected neighbors [n_queries, k] - * @param[in] sample_filter a device filter function that greenlights samples for a given query - */ -template -void search_with_filtering(raft::resources const& handle, - const search_params& params, - const index& index, - raft::device_matrix_view queries, - raft::device_matrix_view neighbors, - raft::device_matrix_view distances, - IvfSampleFilterT sample_filter = IvfSampleFilterT()) -{ - RAFT_EXPECTS( - queries.extent(0) == neighbors.extent(0) && queries.extent(0) == distances.extent(0), - "Number of rows in output neighbors and distances matrices must equal the number of queries."); - - RAFT_EXPECTS(neighbors.extent(1) == distances.extent(1), - "Number of columns in output neighbors and distances matrices must be equal"); - - RAFT_EXPECTS(queries.extent(1) == index.dim(), - "Number of query dimensions should equal number of dimensions in the index."); - - search_with_filtering(handle, - params, - index, - queries.data_handle(), - static_cast(queries.extent(0)), - static_cast(neighbors.extent(1)), - neighbors.data_handle(), - distances.data_handle(), - resource::get_workspace_resource(handle), - sample_filter); -} - -/** - * @brief Search ANN using the constructed index. - * - * See the [ivf_flat::build](#ivf_flat::build) documentation for a usage example. - * - * Note, this function requires a temporary buffer to store intermediate results between cuda kernel - * calls, which may lead to undesirable allocations and slowdown. To alleviate the problem, you can - * pass a pool memory resource or a large enough pre-allocated memory resource to reduce or - * eliminate entirely allocations happening within `search`: - * @code{.cpp} - * ... - * // use default search parameters - * ivf_flat::search_params search_params; - * // Use the same allocator across multiple searches to reduce the number of - * // cuda memory allocations - * ivf_flat::search(handle, search_params, index, queries1, out_inds1, out_dists1); - * ivf_flat::search(handle, search_params, index, queries2, out_inds2, out_dists2); - * ivf_flat::search(handle, search_params, index, queries3, out_inds3, out_dists3); - * ... - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices - * - * @param[in] handle - * @param[in] params configure the search - * @param[in] index ivf-flat constructed index - * @param[in] queries a device pointer to a row-major matrix [n_queries, index->dim()] - * @param[out] neighbors a device pointer to the indices of the neighbors in the source dataset - * [n_queries, k] - * @param[out] distances a device pointer to the distances to the selected neighbors [n_queries, k] - */ -template -void search(raft::resources const& handle, - const search_params& params, - const index& index, - raft::device_matrix_view queries, - raft::device_matrix_view neighbors, - raft::device_matrix_view distances) -{ - search_with_filtering(handle, - params, - index, - queries, - neighbors, - distances, - cuvs::neighbors::filtering::none_ivf_sample_filter()); -} - -/** @} */ - -} // namespace cuvs::neighbors::ivf_flat diff --git a/cpp/include/cuvs/neighbors/ivf_flat.cuh b/cpp/include/cuvs/neighbors/ivf_flat.cuh deleted file mode 100644 index 8fd9628a4..000000000 --- a/cpp/include/cuvs/neighbors/ivf_flat.cuh +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#ifndef RAFT_EXPLICIT_INSTANTIATE_ONLY -#include "ivf_flat-inl.cuh" -#endif - -#ifdef RAFT_COMPILED -#include "ivf_flat-ext.cuh" -#endif diff --git a/cpp/include/cuvs/neighbors/ivf_flat_codepacker.hpp b/cpp/include/cuvs/neighbors/ivf_flat_codepacker.hpp deleted file mode 100644 index 9f1b43380..000000000 --- a/cpp/include/cuvs/neighbors/ivf_flat_codepacker.hpp +++ /dev/null @@ -1,90 +0,0 @@ -/* - * Copyright (c) 2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include -#include -#include -#include -#include - -namespace cuvs::neighbors::ivf_flat::codepacker { - -/** - * Write one flat code into a block by the given offset. The offset indicates the id of the record - * in the list. This function interleaves the code and is intended to later copy the interleaved - * codes over to the IVF list on device. NB: no memory allocation happens here; the block must fit - * the record (offset + 1). - * - * @tparam T - * - * @param[in] flat_code input flat code - * @param[out] block block of memory to write interleaved codes to - * @param[in] dim dimension of the flat code - * @param[in] veclen size of interleaved data chunks - * @param[in] offset how many records to skip before writing the data into the list - */ -template -_RAFT_HOST_DEVICE void pack_1( - const T* flat_code, T* block, uint32_t dim, uint32_t veclen, uint32_t offset) -{ - // The data is written in interleaved groups of `index::kGroupSize` vectors - using interleaved_group = neighbors::detail::div_utils; - - // Interleave dimensions of the source vector while recording it. - // NB: such `veclen` is selected, that `dim % veclen == 0` - auto group_offset = interleaved_group::roundDown(offset); - auto ingroup_id = interleaved_group::mod(offset) * veclen; - - for (uint32_t l = 0; l < dim; l += veclen) { - for (uint32_t j = 0; j < veclen; j++) { - block[group_offset * dim + l * kIndexGroupSize + ingroup_id + j] = flat_code[l + j]; - } - } -} - -/** - * Unpack 1 record of a single list (cluster) in the index to fetch the flat code. The offset - * indicates the id of the record. This function fetches one flat code from an interleaved code. - * - * @tparam T - * - * @param[in] block interleaved block. The block can be thought of as the whole inverted list in - * interleaved format. - * @param[out] flat_code output flat code - * @param[in] dim dimension of the flat code - * @param[in] veclen size of interleaved data chunks - * @param[in] offset fetch the flat code by the given offset - */ -template -_RAFT_HOST_DEVICE void unpack_1( - const T* block, T* flat_code, uint32_t dim, uint32_t veclen, uint32_t offset) -{ - // The data is written in interleaved groups of `index::kGroupSize` vectors - using interleaved_group = neighbors::detail::div_utils; - - // NB: such `veclen` is selected, that `dim % veclen == 0` - auto group_offset = interleaved_group::roundDown(offset); - auto ingroup_id = interleaved_group::mod(offset) * veclen; - - for (uint32_t l = 0; l < dim; l += veclen) { - for (uint32_t j = 0; j < veclen; j++) { - flat_code[l + j] = block[group_offset * dim + l * kIndexGroupSize + ingroup_id + j]; - } - } -} -} // namespace cuvs::neighbors::ivf_flat::codepacker \ No newline at end of file diff --git a/cpp/include/cuvs/neighbors/ivf_flat_helpers.cuh b/cpp/include/cuvs/neighbors/ivf_flat_helpers.cuh deleted file mode 100644 index cca83cea0..000000000 --- a/cpp/include/cuvs/neighbors/ivf_flat_helpers.cuh +++ /dev/null @@ -1,147 +0,0 @@ -/* - * Copyright (c) 2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include -#include - -#include -#include - -#include - -namespace cuvs::neighbors::ivf_flat::helpers { -using namespace cuvs::spatial::knn::detail; // NOLINT -/** - * @defgroup ivf_flat_helpers Helper functions for manipulationg IVF Flat Index - * @{ - */ - -namespace codepacker { - -/** - * Write flat codes into an existing list by the given offset. - * - * NB: no memory allocation happens here; the list must fit the data (offset + n_vec). - * - * Usage example: - * @code{.cpp} - * auto list_data = index.lists()[label]->data.view(); - * // allocate the buffer for the input codes - * auto codes = raft::make_device_matrix(res, n_vec, index.dim()); - * ... prepare n_vecs to pack into the list in codes ... - * // write codes into the list starting from the 42nd position - * ivf_pq::helpers::codepacker::pack( - * res, make_const_mdspan(codes.view()), index.veclen(), 42, list_data); - * @endcode - * - * @tparam T - * @tparam IdxT - * - * @param[in] res - * @param[in] codes flat codes [n_vec, dim] - * @param[in] veclen size of interleaved data chunks - * @param[in] offset how many records to skip before writing the data into the list - * @param[inout] list_data block to write into - */ -template -void pack( - raft::resources const& res, - raft::device_matrix_view codes, - uint32_t veclen, - uint32_t offset, - raft::device_mdspan::list_extents, raft::row_major> - list_data) -{ - cuvs::neighbors::ivf_flat::detail::pack_list_data(res, codes, veclen, offset, list_data); -} - -/** - * @brief Unpack `n_take` consecutive records of a single list (cluster) in the compressed index - * starting at given `offset`. - * - * Usage example: - * @code{.cpp} - * auto list_data = index.lists()[label]->data.view(); - * // allocate the buffer for the output - * uint32_t n_take = 4; - * auto codes = raft::make_device_matrix(res, n_take, index.dim()); - * uint32_t offset = 0; - * // unpack n_take elements from the list - * ivf_pq::helpers::codepacker::unpack(res, list_data, index.veclen(), offset, codes.view()); - * @endcode - * - * @tparam T - * @tparam IdxT - * - * @param[in] res raft resource - * @param[in] list_data block to read from - * @param[in] veclen size of interleaved data chunks - * @param[in] offset - * How many records in the list to skip. - * @param[inout] codes - * the destination buffer [n_take, index.dim()]. - * The length `n_take` defines how many records to unpack, - * it must be <= the list size. - */ -template -void unpack( - raft::resources const& res, - raft::device_mdspan::list_extents, raft::row_major> - list_data, - uint32_t veclen, - uint32_t offset, - raft::device_matrix_view codes) -{ - cuvs::neighbors::ivf_flat::detail::unpack_list_data( - res, list_data, veclen, offset, codes); -} -} // namespace codepacker - -/** - * @brief Public helper API to reset the data and indices ptrs, and the list sizes. Useful for - * externally modifying the index without going through the build stage. The data and indices of the - * IVF lists will be lost. - * - * Usage example: - * @code{.cpp} - * raft::resources res; - * using namespace cuvs::neighbors; - * // use default index parameters - * ivf_flat::index_params index_params; - * // initialize an empty index - * ivf_flat::index index(res, index_params, D); - * // reset the index's state and list sizes - * ivf_flat::helpers::reset_index(res, &index); - * @endcode - * - * @tparam IdxT - * - * @param[in] res raft resource - * @param[inout] index pointer to IVF-PQ index - */ -template -void reset_index(const raft::resources& res, index* index) -{ - auto stream = resource::get_cuda_stream(res); - - utils::memzero(index->list_sizes().data_handle(), index->list_sizes().size(), stream); - utils::memzero(index->data_ptrs().data_handle(), index->data_ptrs().size(), stream); - utils::memzero(index->inds_ptrs().data_handle(), index->inds_ptrs().size(), stream); -} -/** @} */ -} // namespace cuvs::neighbors::ivf_flat::helpers diff --git a/cpp/include/cuvs/neighbors/ivf_flat_serialize.cuh b/cpp/include/cuvs/neighbors/ivf_flat_serialize.cuh deleted file mode 100644 index 37062ea68..000000000 --- a/cpp/include/cuvs/neighbors/ivf_flat_serialize.cuh +++ /dev/null @@ -1,154 +0,0 @@ -/* - * Copyright (c) 2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include "detail/ivf_flat_serialize.cuh" - -namespace cuvs::neighbors::ivf_flat { - -/** - * \defgroup ivf_flat_serialize IVF-Flat Serialize - * @{ - */ - -/** - * Write the index to an output stream - * - * Experimental, both the API and the serialization format are subject to change. - * - * @code{.cpp} - * #include - * - * raft::resources handle; - * - * // create an output stream - * std::ostream os(std::cout.rdbuf()); - * // create an index with `auto index = ivf_flat::build(...);` - * raft::serialize(handle, os, index); - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices - * - * @param[in] handle the raft handle - * @param[in] os output stream - * @param[in] index IVF-Flat index - * - */ -template -void serialize(raft::resources const& handle, std::ostream& os, const index& index) -{ - detail::serialize(handle, os, index); -} - -/** - * Save the index to file. - * - * Experimental, both the API and the serialization format are subject to change. - * - * @code{.cpp} - * #include - * - * raft::resources handle; - * - * // create a string with a filepath - * std::string filename("/path/to/index"); - * // create an index with `auto index = ivf_flat::build(...);` - * raft::serialize(handle, filename, index); - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices - * - * @param[in] handle the raft handle - * @param[in] filename the file name for saving the index - * @param[in] index IVF-Flat index - * - */ -template -void serialize(raft::resources const& handle, - const std::string& filename, - const index& index) -{ - detail::serialize(handle, filename, index); -} - -/** - * Load index from input stream - * - * Experimental, both the API and the serialization format are subject to change. - * - * @code{.cpp} - * #include - * - * raft::resources handle; - * - * // create an input stream - * std::istream is(std::cin.rdbuf()); - * using T = float; // data element type - * using IdxT = int; // type of the index - * auto index = raft::deserialize(handle, is); - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices - * - * @param[in] handle the raft handle - * @param[in] is input stream - * - * @return cuvs::neighbors::ivf_flat::index - */ -template -index deserialize(raft::resources const& handle, std::istream& is) -{ - return detail::deserialize(handle, is); -} - -/** - * Load index from file. - * - * Experimental, both the API and the serialization format are subject to change. - * - * @code{.cpp} - * #include - * - * raft::resources handle; - * - * // create a string with a filepath - * std::string filename("/path/to/index"); - * using T = float; // data element type - * using IdxT = int; // type of the index - * auto index = raft::deserialize(handle, filename); - * @endcode - * - * @tparam T data element type - * @tparam IdxT type of the indices - * - * @param[in] handle the raft handle - * @param[in] filename the name of the file that stores the index - * - * @return cuvs::neighbors::ivf_flat::index - */ -template -index deserialize(raft::resources const& handle, const std::string& filename) -{ - return detail::deserialize(handle, filename); -} - -/**@}*/ - -} // namespace cuvs::neighbors::ivf_flat diff --git a/cpp/include/cuvs/neighbors/ivf_flat_types.hpp b/cpp/include/cuvs/neighbors/ivf_flat_types.hpp deleted file mode 100644 index 28023f474..000000000 --- a/cpp/include/cuvs/neighbors/ivf_flat_types.hpp +++ /dev/null @@ -1,406 +0,0 @@ -/* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include "ann_types.hpp" -#include - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include // std::max -#include -#include -#include - -namespace cuvs::neighbors::ivf_flat { -/** - * @addtogroup ivf_flat - * @{ - */ - -/** Size of the interleaved group (see `index::data` description). */ -constexpr static uint32_t kIndexGroupSize = 32; - -struct index_params : ann::index_params { - /** The number of inverted lists (clusters) */ - uint32_t n_lists = 1024; - /** The number of iterations searching for kmeans centers (index building). */ - uint32_t kmeans_n_iters = 20; - /** The fraction of data to use during iterative kmeans building. */ - double kmeans_trainset_fraction = 0.5; - /** - * By default (adaptive_centers = false), the cluster centers are trained in `ivf_flat::build`, - * and never modified in `ivf_flat::extend`. As a result, you may need to retrain the index - * from scratch after invoking (`ivf_flat::extend`) a few times with new data, the distribution of - * which is no longer representative of the original training set. - * - * The alternative behavior (adaptive_centers = true) is to update the cluster centers for new - * data when it is added. In this case, `index.centers()` are always exactly the centroids of the - * data in the corresponding clusters. The drawback of this behavior is that the centroids depend - * on the order of adding new data (through the classification of the added data); that is, - * `index.centers()` "drift" together with the changing distribution of the newly added data. - */ - bool adaptive_centers = false; - /** - * By default, the algorithm allocates more space than necessary for individual clusters - * (`list_data`). This allows to amortize the cost of memory allocation and reduce the number of - * data copies during repeated calls to `extend` (extending the database). - * - * The alternative is the conservative allocation behavior; when enabled, the algorithm always - * allocates the minimum amount of memory required to store the given number of records. Set this - * flag to `true` if you prefer to use as little GPU memory for the database as possible. - */ - bool conservative_memory_allocation = false; -}; - -struct search_params : ann::search_params { - /** The number of clusters to search. */ - uint32_t n_probes = 20; -}; - -static_assert(std::is_aggregate_v); -static_assert(std::is_aggregate_v); - -template -struct list_spec { - using value_type = ValueT; - using list_extents = raft::matrix_extent; - using index_type = IdxT; - - SizeT align_max; - SizeT align_min; - uint32_t dim; - - constexpr list_spec(uint32_t dim, bool conservative_memory_allocation) - : dim(dim), - align_min(kIndexGroupSize), - align_max(conservative_memory_allocation ? kIndexGroupSize : 1024) - { - } - - // Allow casting between different size-types (for safer size and offset calculations) - template - constexpr explicit list_spec(const list_spec& other_spec) - : dim{other_spec.dim}, align_min{other_spec.align_min}, align_max{other_spec.align_max} - { - } - - /** Determine the extents of an array enough to hold a given amount of data. */ - constexpr auto make_list_extents(SizeT n_rows) const -> list_extents - { - return make_extents(n_rows, dim); - } -}; - -template -using list_data = ivf::list; - -/** - * @brief IVF-flat index. - * - * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset - * - */ -template -struct index : ann::index { - static_assert(!raft::is_narrowing_v, - "IdxT must be able to represent all values of uint32_t"); - - public: - /** - * Vectorized load/store size in elements, determines the size of interleaved data chunks. - * - * TODO: in theory, we can lift this to the template parameter and keep it at hardware maximum - * possible value by padding the `dim` of the data https://github.com/rapidsai/raft/issues/711 - */ - [[nodiscard]] constexpr inline auto veclen() const noexcept -> uint32_t { return veclen_; } - /** Distance metric used for clustering. */ - [[nodiscard]] constexpr inline auto metric() const noexcept -> cuvs::distance::DistanceType - { - return metric_; - } - /** Whether `centers()` change upon extending the index (ivf_pq::extend). */ - [[nodiscard]] constexpr inline auto adaptive_centers() const noexcept -> bool - { - return adaptive_centers_; - } - /** - * Inverted list data [size, dim]. - * - * The data consists of the dataset rows, grouped by their labels (into clusters/lists). - * Within each list (cluster), the data is grouped into blocks of `kIndexGroupSize` interleaved - * vectors. Note, the total index length is slightly larger than the source dataset length, - * because each cluster is padded by `kIndexGroupSize` elements. - * - * Interleaving pattern: - * within groups of `kIndexGroupSize` rows, the data is interleaved with the block size equal to - * `veclen * sizeof(T)`. That is, a chunk of `veclen` consecutive components of one row is - * followed by a chunk of the same size of the next row, and so on. - * - * __Example__: veclen = 2, dim = 6, kIndexGroupSize = 32, list_size = 31 - * - * x[ 0, 0], x[ 0, 1], x[ 1, 0], x[ 1, 1], ... x[14, 0], x[14, 1], x[15, 0], x[15, 1], - * x[16, 0], x[16, 1], x[17, 0], x[17, 1], ... x[30, 0], x[30, 1], - , - , - * x[ 0, 2], x[ 0, 3], x[ 1, 2], x[ 1, 3], ... x[14, 2], x[14, 3], x[15, 2], x[15, 3], - * x[16, 2], x[16, 3], x[17, 2], x[17, 3], ... x[30, 2], x[30, 3], - , - , - * x[ 0, 4], x[ 0, 5], x[ 1, 4], x[ 1, 5], ... x[14, 4], x[14, 5], x[15, 4], x[15, 5], - * x[16, 4], x[16, 5], x[17, 4], x[17, 5], ... x[30, 4], x[30, 5], - , - , - * - */ - /** Sizes of the lists (clusters) [n_lists] - * NB: This may differ from the actual list size if the shared lists have been extended by another - * index - */ - inline auto list_sizes() noexcept -> raft::device_vector_view - { - return list_sizes_.view(); - } - [[nodiscard]] inline auto list_sizes() const noexcept - -> raft::device_vector_view - { - return list_sizes_.view(); - } - - /** k-means cluster centers corresponding to the lists [n_lists, dim] */ - inline auto centers() noexcept -> raft::device_matrix_view - { - return centers_.view(); - } - [[nodiscard]] inline auto centers() const noexcept - -> raft::device_matrix_view - { - return centers_.view(); - } - - /** - * (Optional) Precomputed norms of the `centers` w.r.t. the chosen distance metric [n_lists]. - * - * NB: this may be empty if the index is empty or if the metric does not require the center norms - * calculation. - */ - inline auto center_norms() noexcept -> std::optional> - { - if (center_norms_.has_value()) { - return std::make_optional>(center_norms_->view()); - } else { - return std::nullopt; - } - } - [[nodiscard]] inline auto center_norms() const noexcept - -> std::optional> - { - if (center_norms_.has_value()) { - return std::make_optional>( - center_norms_->view()); - } else { - return std::nullopt; - } - } - - /** Total length of the index. */ - [[nodiscard]] constexpr inline auto size() const noexcept -> IdxT { return total_size_; } - /** Dimensionality of the data. */ - [[nodiscard]] constexpr inline auto dim() const noexcept -> uint32_t - { - return centers_.extent(1); - } - /** Number of clusters/inverted lists. */ - [[nodiscard]] constexpr inline auto n_lists() const noexcept -> uint32_t { return lists_.size(); } - - // Don't allow copying the index for performance reasons (try avoiding copying data) - index(const index&) = delete; - index(index&&) = default; - auto operator=(const index&) -> index& = delete; - auto operator=(index&&) -> index& = default; - ~index() = default; - - /** Construct an empty index. It needs to be trained and then populated. */ - index(raft::resources const& res, - cuvs::distance::DistanceType metric, - uint32_t n_lists, - bool adaptive_centers, - bool conservative_memory_allocation, - uint32_t dim) - : ann::index(), - veclen_(calculate_veclen(dim)), - metric_(metric), - adaptive_centers_(adaptive_centers), - conservative_memory_allocation_{conservative_memory_allocation}, - centers_(raft::make_device_matrix(res, n_lists, dim)), - center_norms_(std::nullopt), - lists_{n_lists}, - list_sizes_{raft::make_device_vector(res, n_lists)}, - data_ptrs_{raft::make_device_vector(res, n_lists)}, - inds_ptrs_{raft::make_device_vector(res, n_lists)}, - total_size_{0} - { - check_consistency(); - } - - /** Construct an empty index. It needs to be trained and then populated. */ - index(raft::resources const& res, const index_params& params, uint32_t dim) - : index(res, - params.metric, - params.n_lists, - params.adaptive_centers, - params.conservative_memory_allocation, - dim) - { - } - - /** Pointers to the inverted lists (clusters) data [n_lists]. */ - inline auto data_ptrs() noexcept -> raft::device_vector_view - { - return data_ptrs_.view(); - } - [[nodiscard]] inline auto data_ptrs() const noexcept - -> raft::device_vector_view - { - return data_ptrs_.view(); - } - - /** Pointers to the inverted lists (clusters) indices [n_lists]. */ - inline auto inds_ptrs() noexcept -> raft::device_vector_view - { - return inds_ptrs_.view(); - } - [[nodiscard]] inline auto inds_ptrs() const noexcept - -> raft::device_vector_view - { - return inds_ptrs_.view(); - } - /** - * Whether to use convervative memory allocation when extending the list (cluster) data - * (see index_params.conservative_memory_allocation). - */ - [[nodiscard]] constexpr inline auto conservative_memory_allocation() const noexcept -> bool - { - return conservative_memory_allocation_; - } - - /** - * Update the state of the dependent index members. - */ - void recompute_internal_state(raft::resources const& res) - { - auto stream = resource::get_cuda_stream(res); - - // Actualize the list pointers - auto this_lists = lists(); - auto this_data_ptrs = data_ptrs(); - auto this_inds_ptrs = inds_ptrs(); - for (uint32_t label = 0; label < this_lists.size(); label++) { - auto& list = this_lists[label]; - const auto data_ptr = list ? list->data.data_handle() : nullptr; - const auto inds_ptr = list ? list->indices.data_handle() : nullptr; - copy(&this_data_ptrs(label), &data_ptr, 1, stream); - copy(&this_inds_ptrs(label), &inds_ptr, 1, stream); - } - auto this_list_sizes = list_sizes().data_handle(); - total_size_ = thrust::reduce(raft::resource::get_thrust_policy(res), - this_list_sizes, - this_list_sizes + this_lists.size(), - 0, - raft::add_op{}); - check_consistency(); - } - - void allocate_center_norms(raft::resources const& res) - { - switch (metric_) { - case cuvs::distance::DistanceType::L2Expanded: - case cuvs::distance::DistanceType::L2SqrtExpanded: - case cuvs::distance::DistanceType::L2Unexpanded: - case cuvs::distance::DistanceType::L2SqrtUnexpanded: - center_norms_ = raft::make_device_vector(res, n_lists()); - break; - default: center_norms_ = std::nullopt; - } - } - - /** Lists' data and indices. */ - inline auto lists() noexcept -> std::vector>>& - { - return lists_; - } - [[nodiscard]] inline auto lists() const noexcept - -> const std::vector>>& - { - return lists_; - } - - private: - /** - * TODO: in theory, we can lift this to the template parameter and keep it at hardware maximum - * possible value by padding the `dim` of the data https://github.com/rapidsai/raft/issues/711 - */ - uint32_t veclen_; - cuvs::distance::DistanceType metric_; - bool adaptive_centers_; - bool conservative_memory_allocation_; - std::vector>> lists_; - raft::device_vector list_sizes_; - raft::device_matrix centers_; - std::optional> center_norms_; - - // Computed members - raft::device_vector data_ptrs_; - raft::device_vector inds_ptrs_; - IdxT total_size_; - - /** Throw an error if the index content is inconsistent. */ - void check_consistency() - { - auto n_lists = lists_.size(); - RAFT_EXPECTS(dim() % veclen_ == 0, "dimensionality is not a multiple of the veclen"); - RAFT_EXPECTS(list_sizes_.extent(0) == n_lists, "inconsistent list size"); - RAFT_EXPECTS(data_ptrs_.extent(0) == n_lists, "inconsistent list size"); - RAFT_EXPECTS(inds_ptrs_.extent(0) == n_lists, "inconsistent list size"); - RAFT_EXPECTS( // - (centers_.extent(0) == list_sizes_.extent(0)) && // - (!center_norms_.has_value() || centers_.extent(0) == center_norms_->extent(0)), - "inconsistent number of lists (clusters)"); - } - - static auto calculate_veclen(uint32_t dim) -> uint32_t - { - // TODO: consider padding the dimensions and fixing veclen to its maximum possible value as a - // template parameter (https://github.com/rapidsai/raft/issues/711) - - // NOTE: keep this consistent with the select_interleaved_scan_kernel logic - // in detail/ivf_flat_interleaved_scan-inl.cuh. - uint32_t veclen = std::max(1, 16 / sizeof(T)); - if (dim % veclen != 0) { veclen = 1; } - return veclen; - } -}; - -/** @} */ - -} // namespace cuvs::neighbors::ivf_flat diff --git a/cpp/include/cuvs/neighbors/ivf_list.hpp b/cpp/include/cuvs/neighbors/ivf_list.hpp deleted file mode 100644 index c395980de..000000000 --- a/cpp/include/cuvs/neighbors/ivf_list.hpp +++ /dev/null @@ -1,196 +0,0 @@ -/* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include -#include -#include - -#include -#include -#include -#include -#include -#include -#include - -#include - -#include -#include -#include - -namespace cuvs::neighbors::ivf { - -/** The data for a single IVF list. */ -template