From d1b53b1b6218bd54df2bbe0eae606fc0e381d90c Mon Sep 17 00:00:00 2001 From: Tamas Bela Feher Date: Sat, 22 Jul 2023 00:09:13 +0200 Subject: [PATCH] Separate cagra index type from internal idx type --- cpp/bench/ann/src/raft/raft_cagra_wrapper.h | 16 ++--- cpp/bench/prims/neighbors/cagra_bench.cuh | 20 +++---- cpp/include/raft/neighbors/cagra.cuh | 46 +++++++-------- cpp/include/raft/neighbors/cagra_types.hpp | 58 +++++++++---------- .../neighbors/detail/cagra/cagra_build.cuh | 11 ++-- .../neighbors/detail/cagra/cagra_search.cuh | 29 +++++----- .../detail/cagra/cagra_serialize.cuh | 6 +- .../neighbors/detail/cagra/graph_core.cuh | 34 ++++++----- .../detail/cagra/search_multi_cta.cuh | 4 +- .../cagra/search_multi_cta_kernel-ext.cuh | 8 +-- .../cagra/search_multi_cta_kernel-inl.cuh | 4 +- .../detail/cagra/search_multi_kernel.cuh | 4 +- .../neighbors/detail/cagra/search_plan.cuh | 4 +- .../detail/cagra/search_single_cta.cuh | 4 +- .../cagra/search_single_cta_kernel-ext.cuh | 8 +-- .../cagra/search_single_cta_kernel-inl.cuh | 4 +- .../cagra/search_multi_cta_00_generate.py | 11 ++-- ...arch_multi_cta_float_uint32_dim1024_t32.cu | 4 +- ...search_multi_cta_float_uint32_dim128_t8.cu | 4 +- ...earch_multi_cta_float_uint32_dim256_t16.cu | 4 +- ...earch_multi_cta_float_uint32_dim512_t32.cu | 4 +- ...arch_multi_cta_float_uint64_dim1024_t32.cu | 4 +- ...search_multi_cta_float_uint64_dim128_t8.cu | 4 +- ...earch_multi_cta_float_uint64_dim256_t16.cu | 4 +- ...earch_multi_cta_float_uint64_dim512_t32.cu | 4 +- ...earch_multi_cta_int8_uint32_dim1024_t32.cu | 4 +- .../search_multi_cta_int8_uint32_dim128_t8.cu | 4 +- ...search_multi_cta_int8_uint32_dim256_t16.cu | 4 +- ...search_multi_cta_int8_uint32_dim512_t32.cu | 4 +- ...arch_multi_cta_uint8_uint32_dim1024_t32.cu | 4 +- ...search_multi_cta_uint8_uint32_dim128_t8.cu | 4 +- ...earch_multi_cta_uint8_uint32_dim256_t16.cu | 4 +- ...earch_multi_cta_uint8_uint32_dim512_t32.cu | 4 +- .../cagra/search_single_cta_00_generate.py | 4 +- ...rch_single_cta_float_uint32_dim1024_t32.cu | 4 +- ...earch_single_cta_float_uint32_dim128_t8.cu | 4 +- ...arch_single_cta_float_uint32_dim256_t16.cu | 4 +- ...arch_single_cta_float_uint32_dim512_t32.cu | 4 +- ...rch_single_cta_float_uint64_dim1024_t32.cu | 4 +- ...earch_single_cta_float_uint64_dim128_t8.cu | 4 +- ...arch_single_cta_float_uint64_dim256_t16.cu | 4 +- ...arch_single_cta_float_uint64_dim512_t32.cu | 4 +- ...arch_single_cta_int8_uint32_dim1024_t32.cu | 4 +- ...search_single_cta_int8_uint32_dim128_t8.cu | 4 +- ...earch_single_cta_int8_uint32_dim256_t16.cu | 4 +- ...earch_single_cta_int8_uint32_dim512_t32.cu | 4 +- ...rch_single_cta_uint8_uint32_dim1024_t32.cu | 4 +- ...earch_single_cta_uint8_uint32_dim128_t8.cu | 4 +- ...arch_single_cta_uint8_uint32_dim256_t16.cu | 4 +- ...arch_single_cta_uint8_uint32_dim512_t32.cu | 4 +- cpp/test/neighbors/ann_cagra.cuh | 30 +++++----- .../ann_cagra/search_kernel_uint64_t.cuh | 8 +-- 52 files changed, 224 insertions(+), 217 deletions(-) diff --git a/cpp/bench/ann/src/raft/raft_cagra_wrapper.h b/cpp/bench/ann/src/raft/raft_cagra_wrapper.h index 79ae746078..e6490b3f07 100644 --- a/cpp/bench/ann/src/raft/raft_cagra_wrapper.h +++ b/cpp/bench/ann/src/raft/raft_cagra_wrapper.h @@ -71,7 +71,7 @@ class RaftCagra : public ANN { AlgoProperty get_property() const override { AlgoProperty property; - property.dataset_memory_type = MemoryType::Host; + property.dataset_memory_type = MemoryType::HostMmap; property.query_memory_type = MemoryType::Device; property.need_dataset_when_search = true; return property; @@ -104,13 +104,14 @@ RaftCagra::RaftCagra(Metric metric, int dim, const BuildParam& param) template void RaftCagra::build(const T* dataset, size_t nrow, cudaStream_t) { - if (get_property().dataset_memory_type == MemoryType::Host) { - auto dataset_view = raft::make_host_matrix_view(dataset, IdxT(nrow), dimension_); + if (get_property().dataset_memory_type != MemoryType::Device) { + auto dataset_view = + raft::make_host_matrix_view(dataset, IdxT(nrow), dimension_); index_.emplace( raft::neighbors::experimental::cagra::build(handle_, index_params_, dataset_view)); } else { auto dataset_view = - raft::make_device_matrix_view(dataset, IdxT(nrow), dimension_); + raft::make_device_matrix_view(dataset, IdxT(nrow), dimension_); index_.emplace( raft::neighbors::experimental::cagra::build(handle_, index_params_, dataset_view)); } @@ -152,9 +153,10 @@ void RaftCagra::search( neighbors_IdxT = neighbors_storage.data(); } - auto queries_view = raft::make_device_matrix_view(queries, batch_size, dimension_); - auto neighbors_view = raft::make_device_matrix_view(neighbors_IdxT, batch_size, k); - auto distances_view = raft::make_device_matrix_view(distances, batch_size, k); + auto queries_view = + raft::make_device_matrix_view(queries, batch_size, dimension_); + auto neighbors_view = raft::make_device_matrix_view(neighbors_IdxT, batch_size, k); + auto distances_view = raft::make_device_matrix_view(distances, batch_size, k); raft::neighbors::experimental::cagra::search( handle_, search_params_, *index_, queries_view, neighbors_view, distances_view); diff --git a/cpp/bench/prims/neighbors/cagra_bench.cuh b/cpp/bench/prims/neighbors/cagra_bench.cuh index c361dc82dc..19679377d8 100644 --- a/cpp/bench/prims/neighbors/cagra_bench.cuh +++ b/cpp/bench/prims/neighbors/cagra_bench.cuh @@ -47,9 +47,9 @@ struct CagraBench : public fixture { explicit CagraBench(const params& ps) : fixture(true), params_(ps), - queries_(make_device_matrix(handle, ps.n_queries, ps.n_dims)), - dataset_(make_device_matrix(handle, ps.n_samples, ps.n_dims)), - knn_graph_(make_device_matrix(handle, ps.n_samples, ps.degree)) + queries_(make_device_matrix(handle, ps.n_queries, ps.n_dims)), + dataset_(make_device_matrix(handle, ps.n_samples, ps.n_dims)), + knn_graph_(make_device_matrix(handle, ps.n_samples, ps.degree)) { // Generate random dataset and queriees raft::random::RngState state{42}; @@ -87,11 +87,11 @@ struct CagraBench : public fixture { search_params.thread_block_size = params_.block_size; search_params.num_parents = params_.num_parents; - auto indices = make_device_matrix(handle, params_.n_queries, params_.k); - auto distances = make_device_matrix(handle, params_.n_queries, params_.k); - auto ind_v = make_device_matrix_view( + auto indices = make_device_matrix(handle, params_.n_queries, params_.k); + auto distances = make_device_matrix(handle, params_.n_queries, params_.k); + auto ind_v = make_device_matrix_view( indices.data_handle(), params_.n_queries, params_.k); - auto dist_v = make_device_matrix_view( + auto dist_v = make_device_matrix_view( distances.data_handle(), params_.n_queries, params_.k); auto queries_v = make_const_mdspan(queries_.view()); @@ -125,9 +125,9 @@ struct CagraBench : public fixture { private: const params params_; std::optional> index_; - raft::device_matrix queries_; - raft::device_matrix dataset_; - raft::device_matrix knn_graph_; + raft::device_matrix queries_; + raft::device_matrix dataset_; + raft::device_matrix knn_graph_; }; inline const std::vector generate_inputs() diff --git a/cpp/include/raft/neighbors/cagra.cuh b/cpp/include/raft/neighbors/cagra.cuh index 5934f6ef69..c9f2127572 100644 --- a/cpp/include/raft/neighbors/cagra.cuh +++ b/cpp/include/raft/neighbors/cagra.cuh @@ -65,7 +65,7 @@ namespace raft::neighbors::experimental::cagra { * @endcode * * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset + * @tparam IdxT type of the dataset vector indices * * @param[in] res raft resources * @param[in] dataset a matrix view (host or device) to a row-major matrix [n_rows, dim] @@ -76,19 +76,19 @@ namespace raft::neighbors::experimental::cagra { */ template void build_knn_graph(raft::resources const& res, - mdspan, row_major, accessor> dataset, - raft::host_matrix_view knn_graph, + mdspan, row_major, accessor> dataset, + raft::host_matrix_view knn_graph, std::optional refine_rate = std::nullopt, std::optional build_params = std::nullopt, std::optional search_params = std::nullopt) { using internal_IdxT = typename std::make_unsigned::type; - auto knn_graph_internal = make_host_matrix_view( + auto knn_graph_internal = make_host_matrix_view( reinterpret_cast(knn_graph.data_handle()), knn_graph.extent(0), knn_graph.extent(1)); - auto dataset_internal = mdspan, row_major, accessor>( + auto dataset_internal = mdspan, row_major, accessor>( dataset.data_handle(), dataset.extent(0), dataset.extent(1)); detail::build_knn_graph( @@ -119,7 +119,7 @@ void build_knn_graph(raft::resources const& res, * @endcode * * @tparam DataT type of the data in the source dataset - * @tparam IdxT type of the indices in the source dataset + * @tparam IdxT type of the dataset vector indices * * @param[in] res raft resources * @param[in] dataset a matrix view (host or device) to a row-major matrix [n_rows, dim] @@ -133,20 +133,20 @@ template , memory_type::host>> void sort_knn_graph(raft::resources const& res, - mdspan, row_major, d_accessor> dataset, - mdspan, row_major, g_accessor> knn_graph) + mdspan, row_major, d_accessor> dataset, + mdspan, row_major, g_accessor> knn_graph) { using internal_IdxT = typename std::make_unsigned::type; using g_accessor_internal = host_device_accessor, g_accessor::mem_type>; auto knn_graph_internal = - mdspan, row_major, g_accessor_internal>( + mdspan, row_major, g_accessor_internal>( reinterpret_cast(knn_graph.data_handle()), knn_graph.extent(0), knn_graph.extent(1)); - auto dataset_internal = mdspan, row_major, d_accessor>( + auto dataset_internal = mdspan, row_major, d_accessor>( dataset.data_handle(), dataset.extent(0), dataset.extent(1)); detail::graph::sort_knn_graph(res, dataset_internal, knn_graph_internal); @@ -170,12 +170,12 @@ template , memory_type::host>> void optimize(raft::resources const& res, - mdspan, row_major, g_accessor> knn_graph, - raft::host_matrix_view new_graph) + mdspan, row_major, g_accessor> knn_graph, + raft::host_matrix_view new_graph) { using internal_IdxT = typename std::make_unsigned::type; - auto new_graph_internal = raft::make_host_matrix_view( + auto new_graph_internal = raft::make_host_matrix_view( reinterpret_cast(new_graph.data_handle()), new_graph.extent(0), new_graph.extent(1)); @@ -183,7 +183,7 @@ void optimize(raft::resources const& res, using g_accessor_internal = host_device_accessor, memory_type::host>; auto knn_graph_internal = - mdspan, row_major, g_accessor_internal>( + mdspan, row_major, g_accessor_internal>( reinterpret_cast(knn_graph.data_handle()), knn_graph.extent(0), knn_graph.extent(1)); @@ -237,7 +237,7 @@ template , memory_type::host>> index build(raft::resources const& res, const index_params& params, - mdspan, row_major, Accessor> dataset) + mdspan, row_major, Accessor> dataset) { size_t intermediate_degree = params.intermediate_graph_degree; size_t graph_degree = params.graph_degree; @@ -256,11 +256,11 @@ index build(raft::resources const& res, graph_degree = intermediate_degree; } - auto knn_graph = raft::make_host_matrix(dataset.extent(0), intermediate_degree); + auto knn_graph = raft::make_host_matrix(dataset.extent(0), intermediate_degree); build_knn_graph(res, dataset, knn_graph.view()); - auto cagra_graph = raft::make_host_matrix(dataset.extent(0), graph_degree); + auto cagra_graph = raft::make_host_matrix(dataset.extent(0), graph_degree); optimize(res, knn_graph.view(), cagra_graph.view()); @@ -289,9 +289,9 @@ template void search(raft::resources const& res, const search_params& params, const index& idx, - raft::device_matrix_view queries, - raft::device_matrix_view neighbors, - raft::device_matrix_view distances) + raft::device_matrix_view queries, + raft::device_matrix_view neighbors, + raft::device_matrix_view distances) { RAFT_EXPECTS( queries.extent(0) == neighbors.extent(0) && queries.extent(0) == distances.extent(0), @@ -303,13 +303,13 @@ void search(raft::resources const& res, "Number of query dimensions should equal number of dimensions in the index."); using internal_IdxT = typename std::make_unsigned::type; - auto queries_internal = raft::make_device_matrix_view( + auto queries_internal = raft::make_device_matrix_view( queries.data_handle(), queries.extent(0), queries.extent(1)); - auto neighbors_internal = raft::make_device_matrix_view( + auto neighbors_internal = raft::make_device_matrix_view( reinterpret_cast(neighbors.data_handle()), neighbors.extent(0), neighbors.extent(1)); - auto distances_internal = raft::make_device_matrix_view( + auto distances_internal = raft::make_device_matrix_view( distances.data_handle(), distances.extent(0), distances.extent(1)); detail::search_main( diff --git a/cpp/include/raft/neighbors/cagra_types.hpp b/cpp/include/raft/neighbors/cagra_types.hpp index 44375c01f0..c046260cab 100644 --- a/cpp/include/raft/neighbors/cagra_types.hpp +++ b/cpp/include/raft/neighbors/cagra_types.hpp @@ -108,7 +108,7 @@ static_assert(std::is_aggregate_v); * The index stores the dataset and a kNN graph in device memory. * * @tparam T data element type - * @tparam IdxT type of the indices in the source dataset + * @tparam IdxT type of the vector indices (represent dataset.extent(0)) * */ template @@ -124,7 +124,7 @@ struct index : ann::index { return metric_; } - // /** Total length of the index. */ + // /** Total length of the index (number of vectors). */ [[nodiscard]] constexpr inline auto size() const noexcept -> IdxT { return dataset_view_.extent(0); @@ -143,16 +143,14 @@ struct index : ann::index { /** Dataset [size, dim] */ [[nodiscard]] inline auto dataset() const noexcept - -> device_matrix_view + -> device_matrix_view { return dataset_view_; } /** neighborhood graph [size, graph-degree] */ - inline auto graph() noexcept -> device_matrix_view { return graph_view_; } - [[nodiscard]] inline auto graph() const noexcept - -> device_matrix_view + -> device_matrix_view { return graph_view_; } @@ -168,8 +166,8 @@ struct index : ann::index { index(raft::resources const& res) : ann::index(), metric_(raft::distance::DistanceType::L2Expanded), - dataset_(make_device_matrix(res, 0, 0)), - graph_(make_device_matrix(res, 0, 0)) + dataset_(make_device_matrix(res, 0, 0)), + graph_(make_device_matrix(res, 0, 0)) { } @@ -189,7 +187,7 @@ struct index : ann::index { * - Cagra index is normally created by the cagra::build * @code{.cpp} * using namespace raft::neighbors::experimental; - * auto dataset = raft::make_host_matrix(n_rows, n_cols); + * auto dataset = raft::make_host_matrix(n_rows, n_cols); * load_dataset(dataset.view()); * // use default index parameters * cagra::index_params index_params; @@ -198,8 +196,8 @@ struct index : ann::index { * // 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); + * 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 @@ -210,8 +208,8 @@ struct index : ann::index { * @code{.cpp} * using namespace raft::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); + * 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()); @@ -230,12 +228,12 @@ struct index : ann::index { template index(raft::resources const& res, raft::distance::DistanceType metric, - mdspan, row_major, data_accessor> dataset, - mdspan, row_major, graph_accessor> knn_graph) + mdspan, row_major, data_accessor> dataset, + mdspan, row_major, graph_accessor> knn_graph) : ann::index(), metric_(metric), - dataset_(make_device_matrix(res, 0, 0)), - graph_(make_device_matrix(res, 0, 0)) + dataset_(make_device_matrix(res, 0, 0)), + graph_(make_device_matrix(res, 0, 0)) { RAFT_EXPECTS(dataset.extent(0) == knn_graph.extent(0), "Dataset and knn_graph must have equal number of rows"); @@ -252,13 +250,13 @@ struct index : ann::index { * index. */ void update_dataset(raft::resources const& res, - raft::device_matrix_view dataset) + raft::device_matrix_view dataset) { if (dataset.extent(1) % AlignDim::Value != 0) { RAFT_LOG_DEBUG("Creating a padded copy of CAGRA dataset in device memory"); copy_padded(res, dataset); } else { - dataset_view_ = make_device_strided_matrix_view( + dataset_view_ = make_device_strided_matrix_view( dataset.data_handle(), dataset.extent(0), dataset.extent(1), dataset.extent(1)); } } @@ -269,7 +267,7 @@ struct index : ann::index { * 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::host_matrix_view dataset) { RAFT_LOG_DEBUG("Copying CAGRA dataset from host to device"); copy_padded(res, dataset); @@ -282,7 +280,7 @@ struct index : ann::index { * 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::device_matrix_view knn_graph) { graph_view_ = knn_graph; } @@ -293,10 +291,10 @@ struct index : ann::index { * 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::host_matrix_view knn_graph) { RAFT_LOG_DEBUG("Copying CAGRA knn graph from host to device"); - graph_ = make_device_matrix(res, knn_graph.extent(0), knn_graph.extent(1)); + graph_ = make_device_matrix(res, knn_graph.extent(0), knn_graph.extent(1)); raft::copy(graph_.data_handle(), knn_graph.data_handle(), knn_graph.size(), @@ -308,10 +306,10 @@ struct index : ann::index { /** Create a device copy of the dataset, and pad it if necessary. */ template void copy_padded(raft::resources const& res, - mdspan, row_major, data_accessor> dataset) + mdspan, row_major, data_accessor> dataset) { dataset_ = - make_device_matrix(res, dataset.extent(0), AlignDim::roundUp(dataset.extent(1))); + make_device_matrix(res, dataset.extent(0), AlignDim::roundUp(dataset.extent(1))); if (dataset_.extent(1) == dataset.extent(1)) { raft::copy(dataset_.data_handle(), dataset.data_handle(), @@ -330,7 +328,7 @@ struct index : ann::index { cudaMemcpyDefault, resource::get_cuda_stream(res))); } - dataset_view_ = make_device_strided_matrix_view( + dataset_view_ = make_device_strided_matrix_view( dataset_.data_handle(), dataset_.extent(0), dataset.extent(1), dataset_.extent(1)); RAFT_LOG_DEBUG("CAGRA dataset strided matrix view %zux%zu, stride %zu", static_cast(dataset_view_.extent(0)), @@ -339,10 +337,10 @@ struct index : ann::index { } raft::distance::DistanceType metric_; - raft::device_matrix dataset_; - raft::device_matrix graph_; - raft::device_matrix_view dataset_view_; - raft::device_matrix_view graph_view_; + raft::device_matrix dataset_; + raft::device_matrix graph_; + raft::device_matrix_view dataset_view_; + raft::device_matrix_view graph_view_; }; /** @} */ diff --git a/cpp/include/raft/neighbors/detail/cagra/cagra_build.cuh b/cpp/include/raft/neighbors/detail/cagra/cagra_build.cuh index d2bf7bf1ed..4f8323a481 100644 --- a/cpp/include/raft/neighbors/detail/cagra/cagra_build.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/cagra_build.cuh @@ -40,8 +40,8 @@ namespace raft::neighbors::experimental::cagra::detail { template void build_knn_graph(raft::resources const& res, - mdspan, row_major, accessor> dataset, - raft::host_matrix_view knn_graph, + mdspan, row_major, accessor> dataset, + raft::host_matrix_view knn_graph, std::optional refine_rate = std::nullopt, std::optional build_params = std::nullopt, std::optional search_params = std::nullopt) @@ -108,7 +108,6 @@ void build_knn_graph(raft::resources const& res, max_batch_size, search_params->n_probes); - // TODO(tfeher): shall we use uint32_t? auto distances = raft::make_device_matrix(res, max_batch_size, gpu_top_k); auto neighbors = raft::make_device_matrix(res, max_batch_size, gpu_top_k); auto refined_distances = raft::make_device_matrix(res, max_batch_size, top_k); @@ -139,6 +138,8 @@ void build_knn_graph(raft::resources const& res, size_t d_report_offset = dataset.extent(0) / 100; // Report progress in 1% steps. for (const auto& batch : vec_batches) { + // Map int64_t to uint32_t because ivf_pq requires the latter. + // TODO(tfeher): remove this mapping once ivf_pq accepts mdspan with int64_t index type auto queries_view = raft::make_device_matrix_view( batch.data(), batch.size(), batch.row_width()); auto neighbors_view = make_device_matrix_view( @@ -147,7 +148,6 @@ void build_knn_graph(raft::resources const& res, distances.data_handle(), batch.size(), distances.extent(1)); ivf_pq::search(res, *search_params, index, queries_view, neighbors_view, distances_view); - if constexpr (is_host_mdspan_v) { raft::copy(neighbors_host.data_handle(), neighbors.data_handle(), @@ -167,7 +167,7 @@ void build_knn_graph(raft::resources const& res, refined_distances_host.data_handle(), batch.size(), top_k); resource::sync_stream(res); - raft::neighbors::detail::refine_host( // res, + raft::neighbors::detail::refine_host( dataset, queries_host_view, neighbors_host_view, @@ -234,6 +234,7 @@ void build_knn_graph(raft::resources const& res, } first = false; } + if (!first) RAFT_LOG_DEBUG("# Finished building kNN graph"); } diff --git a/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh b/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh index 1561a3bb8d..e7688db1af 100644 --- a/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh @@ -38,7 +38,9 @@ namespace raft::neighbors::experimental::cagra::detail { * See the [build](#build) documentation for a usage example. * * @tparam T data element type - * @tparam IdxT type of the indices + * @tparam IdxT type of database vector indices + * @tparam internal_IdxT during search we map IdxT to internal_IdxT, this way we do not need + * separate kernels for int/uint. * * @param[in] handle * @param[in] params configure the search @@ -54,9 +56,9 @@ template & index, - raft::device_matrix_view queries, - raft::device_matrix_view neighbors, - raft::device_matrix_view distances) + raft::device_matrix_view queries, + raft::device_matrix_view neighbors, + raft::device_matrix_view distances) { RAFT_LOG_DEBUG("# dataset size = %lu, dim = %lu\n", static_cast(index.dataset().extent(0)), @@ -92,16 +94,15 @@ void search_main(raft::resources const& res, : nullptr; uint32_t* _num_executed_iterations = nullptr; - auto dataset_internal = make_device_strided_matrix_view( - index.dataset().data_handle(), - index.dataset().extent(0), - index.dataset().extent(1), - index.dataset().stride(0)); - auto graph_internal = - raft::make_device_matrix_view( - reinterpret_cast(index.graph().data_handle()), - index.graph().extent(0), - index.graph().extent(1)); + auto dataset_internal = + make_device_strided_matrix_view(index.dataset().data_handle(), + index.dataset().extent(0), + index.dataset().extent(1), + index.dataset().stride(0)); + auto graph_internal = raft::make_device_matrix_view( + reinterpret_cast(index.graph().data_handle()), + index.graph().extent(0), + index.graph().extent(1)); (*plan)(res, dataset_internal, diff --git a/cpp/include/raft/neighbors/detail/cagra/cagra_serialize.cuh b/cpp/include/raft/neighbors/detail/cagra/cagra_serialize.cuh index 7f708506a5..1b95daf431 100644 --- a/cpp/include/raft/neighbors/detail/cagra/cagra_serialize.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/cagra_serialize.cuh @@ -62,7 +62,7 @@ void serialize(raft::resources const& res, std::ostream& os, const index(dataset.extent(0), dataset.extent(1)); + auto host_dataset = make_host_matrix(dataset.extent(0), dataset.extent(1)); RAFT_CUDA_TRY(cudaMemcpy2DAsync(host_dataset.data_handle(), sizeof(T) * host_dataset.extent(1), dataset.data_handle(), @@ -111,8 +111,8 @@ auto deserialize(raft::resources const& res, std::istream& is) -> index auto graph_degree = deserialize_scalar(res, is); auto metric = deserialize_scalar(res, is); - auto dataset = raft::make_host_matrix(n_rows, dim); - auto graph = raft::make_host_matrix(n_rows, graph_degree); + auto dataset = raft::make_host_matrix(n_rows, dim); + auto graph = raft::make_host_matrix(n_rows, graph_degree); deserialize_mdspan(res, is, dataset.view()); deserialize_mdspan(res, is, graph.view()); diff --git a/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh b/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh index d915634df9..f66d85a6b8 100644 --- a/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh @@ -231,8 +231,8 @@ template , memory_type::host>> void sort_knn_graph(raft::resources const& res, - mdspan, row_major, d_accessor> dataset, - mdspan, row_major, g_accessor> knn_graph) + mdspan, row_major, d_accessor> dataset, + mdspan, row_major, g_accessor> knn_graph) { RAFT_EXPECTS(dataset.extent(0) == knn_graph.extent(0), "dataset size is expected to have the same number of graph index size"); @@ -252,7 +252,7 @@ void sort_knn_graph(raft::resources const& res, const double time_sort_start = cur_time(); RAFT_LOG_DEBUG("# Sorting kNN Graph on GPUs "); - auto d_dataset = raft::make_device_matrix(res, dataset_size, dataset_dim); + auto d_dataset = raft::make_device_matrix(res, dataset_size, dataset_dim); raft::copy(d_dataset.data_handle(), dataset_ptr, dataset_size * dataset_dim, @@ -318,8 +318,8 @@ template , memory_type::host>> void optimize(raft::resources const& res, - mdspan, row_major, g_accessor> knn_graph, - raft::host_matrix_view new_graph) + mdspan, row_major, g_accessor> knn_graph, + raft::host_matrix_view new_graph) { RAFT_LOG_DEBUG( "# Pruning kNN graph (size=%lu, degree=%lu)\n", knn_graph.extent(0), knn_graph.extent(1)); @@ -334,23 +334,24 @@ void optimize(raft::resources const& res, auto output_graph_ptr = new_graph.data_handle(); const IdxT graph_size = new_graph.extent(0); - auto pruned_graph = raft::make_host_matrix(graph_size, output_graph_degree); + auto pruned_graph = raft::make_host_matrix(graph_size, output_graph_degree); { // // Prune kNN graph // - auto d_input_graph = raft::make_device_matrix(res, graph_size, input_graph_degree); + auto d_input_graph = + raft::make_device_matrix(res, graph_size, input_graph_degree); - auto detour_count = raft::make_host_matrix(graph_size, input_graph_degree); + auto detour_count = raft::make_host_matrix(graph_size, input_graph_degree); auto d_detour_count = - raft::make_device_matrix(res, graph_size, input_graph_degree); + raft::make_device_matrix(res, graph_size, input_graph_degree); RAFT_CUDA_TRY(cudaMemsetAsync(d_detour_count.data_handle(), 0xff, graph_size * input_graph_degree * sizeof(uint8_t), resource::get_cuda_stream(res))); - auto d_num_no_detour_edges = raft::make_device_vector(res, graph_size); + auto d_num_no_detour_edges = raft::make_device_vector(res, graph_size); RAFT_CUDA_TRY(cudaMemsetAsync(d_num_no_detour_edges.data_handle(), 0x00, graph_size * sizeof(uint32_t), @@ -468,8 +469,8 @@ void optimize(raft::resources const& res, (double)num_full / graph_size * 100); } - auto rev_graph = raft::make_host_matrix(graph_size, output_graph_degree); - auto rev_graph_count = raft::make_host_vector(graph_size); + auto rev_graph = raft::make_host_matrix(graph_size, output_graph_degree); + auto rev_graph_count = raft::make_host_vector(graph_size); { // @@ -477,20 +478,21 @@ void optimize(raft::resources const& res, // const double time_make_start = cur_time(); - auto d_rev_graph = raft::make_device_matrix(res, graph_size, output_graph_degree); + auto d_rev_graph = + raft::make_device_matrix(res, graph_size, output_graph_degree); RAFT_CUDA_TRY(cudaMemsetAsync(d_rev_graph.data_handle(), 0xff, graph_size * output_graph_degree * sizeof(IdxT), resource::get_cuda_stream(res))); - auto d_rev_graph_count = raft::make_device_vector(res, graph_size); + auto d_rev_graph_count = raft::make_device_vector(res, graph_size); RAFT_CUDA_TRY(cudaMemsetAsync(d_rev_graph_count.data_handle(), 0x00, graph_size * sizeof(uint32_t), resource::get_cuda_stream(res))); - auto dest_nodes = raft::make_host_vector(graph_size); - auto d_dest_nodes = raft::make_device_vector(res, graph_size); + auto dest_nodes = raft::make_host_vector(graph_size); + auto d_dest_nodes = raft::make_device_vector(res, graph_size); for (uint64_t k = 0; k < output_graph_degree; k++) { #pragma omp parallel for diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta.cuh index bf6a32eac8..b8a05239be 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta.cuh @@ -178,8 +178,8 @@ struct search : public search_plan_impl { ~search() {} void operator()(raft::resources const& res, - raft::device_matrix_view dataset, - raft::device_matrix_view graph, + raft::device_matrix_view dataset, + raft::device_matrix_view graph, INDEX_T* const topk_indices_ptr, // [num_queries, topk] DISTANCE_T* const topk_distances_ptr, // [num_queries, topk] const DATA_T* const queries_ptr, // [num_queries, dataset_dim] diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh index 3ccd73d92c..8f5cfa08ce 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh @@ -27,8 +27,8 @@ template -void select_and_run(raft::device_matrix_view dataset, - raft::device_matrix_view graph, +void select_and_run(raft::device_matrix_view dataset, + raft::device_matrix_view graph, INDEX_T* const topk_indices_ptr, DISTANCE_T* const topk_distances_ptr, const DATA_T* const queries_ptr, @@ -54,8 +54,8 @@ void select_and_run(raft::device_matrix_view( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh index 43e3e83f59..ad3012a9ae 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh @@ -452,8 +452,8 @@ template void select_and_run( // raft::resources const& res, - raft::device_matrix_view dataset, - raft::device_matrix_view graph, + raft::device_matrix_view dataset, + raft::device_matrix_view graph, INDEX_T* const topk_indices_ptr, // [num_queries, topk] DISTANCE_T* const topk_distances_ptr, // [num_queries, topk] const DATA_T* const queries_ptr, // [num_queries, dataset_dim] diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh index 033022aea1..4e1d7955f0 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh @@ -594,8 +594,8 @@ struct search : search_plan_impl { ~search() {} void operator()(raft::resources const& res, - raft::device_matrix_view dataset, - raft::device_matrix_view graph, + raft::device_matrix_view dataset, + raft::device_matrix_view graph, INDEX_T* const topk_indices_ptr, // [num_queries, topk] DISTANCE_T* const topk_distances_ptr, // [num_queries, topk] const DATA_T* const queries_ptr, // [num_queries, dataset_dim] diff --git a/cpp/include/raft/neighbors/detail/cagra/search_plan.cuh b/cpp/include/raft/neighbors/detail/cagra/search_plan.cuh index cbffd93caf..39e71733ce 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_plan.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_plan.cuh @@ -105,8 +105,8 @@ struct search_plan_impl : public search_plan_impl_base { virtual ~search_plan_impl() {} virtual void operator()(raft::resources const& res, - raft::device_matrix_view dataset, - raft::device_matrix_view graph, + raft::device_matrix_view dataset, + raft::device_matrix_view graph, INDEX_T* const result_indices_ptr, // [num_queries, topk] DISTANCE_T* const result_distances_ptr, // [num_queries, topk] const DATA_T* const queries_ptr, // [num_queries, dataset_dim] diff --git a/cpp/include/raft/neighbors/detail/cagra/search_single_cta.cuh b/cpp/include/raft/neighbors/detail/cagra/search_single_cta.cuh index bad2039f8c..a7576bb82b 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_single_cta.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_single_cta.cuh @@ -226,8 +226,8 @@ struct search : search_plan_impl { } void operator()(raft::resources const& res, - raft::device_matrix_view dataset, - raft::device_matrix_view graph, + raft::device_matrix_view dataset, + raft::device_matrix_view graph, INDEX_T* const result_indices_ptr, // [num_queries, topk] DISTANCE_T* const result_distances_ptr, // [num_queries, topk] const DATA_T* const queries_ptr, // [num_queries, dataset_dim] diff --git a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh index b0130e45d4..a8715620e3 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh @@ -27,8 +27,8 @@ template void select_and_run( // raft::resources const& res, - raft::device_matrix_view dataset, - raft::device_matrix_view graph, + raft::device_matrix_view dataset, + raft::device_matrix_view graph, INDEX_T* const topk_indices_ptr, // [num_queries, topk] DISTANCE_T* const topk_distances_ptr, // [num_queries, topk] const DATA_T* const queries_ptr, // [num_queries, dataset_dim] @@ -57,8 +57,8 @@ void select_and_run( // raft::resources const& res, #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ extern template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh index ca2166ab8d..eeef4496ff 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh @@ -828,8 +828,8 @@ template void select_and_run( // raft::resources const& res, - raft::device_matrix_view dataset, - raft::device_matrix_view graph, + raft::device_matrix_view dataset, + raft::device_matrix_view graph, INDEX_T* const topk_indices_ptr, // [num_queries, topk] DISTANCE_T* const topk_distances_ptr, // [num_queries, topk] const DATA_T* const queries_ptr, // [num_queries, dataset_dim] diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_00_generate.py b/cpp/src/neighbors/detail/cagra/search_multi_cta_00_generate.py index 170c57c521..3203893ded 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_00_generate.py +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_00_generate.py @@ -44,8 +44,8 @@ #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \\ template void select_and_run( \\ - raft::device_matrix_view dataset, \\ - raft::device_matrix_view graph, \\ + raft::device_matrix_view dataset, \\ + raft::device_matrix_view graph, \\ INDEX_T* const topk_indices_ptr, \\ DISTANCE_T* const topk_distances_ptr, \\ const DATA_T* const queries_ptr, \\ @@ -81,12 +81,15 @@ # mxelem = [64, 128, 256] load_types = ["uint4"] search_types = dict( - float_uint32=("float", "uint32_t", "float"), # data_t, idx_t, distance_t + float_uint32=( + "float", + "uint32_t", + "float", + ), # data_t, vec_idx_t, distance_t int8_uint32=("int8_t", "uint32_t", "float"), uint8_uint32=("uint8_t", "uint32_t", "float"), float_uint64=("float", "uint64_t", "float"), ) - # knn for type_path, (data_t, idx_t, distance_t) in search_types.items(): for (mxdim, team) in mxdim_team: diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu index 207028dcec..3593f61f34 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu index 4a5c0f106b..adce4ce92f 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu index 93a9f41881..29d29c726f 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu index fb321b2cf7..22bd83e801 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim1024_t32.cu index e73698460d..9dfb2b6d6a 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim1024_t32.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim128_t8.cu index e51fdcbc62..0d83a04a6e 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim128_t8.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim256_t16.cu index caa45b5395..23831c9407 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim256_t16.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim512_t32.cu index 67e54f0937..a090ff7d7a 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim512_t32.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu index 2e929eb4f0..d41ed6dfe8 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu index d3e2e78250..572061a3d7 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu index 802edafdf2..dc77e68171 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu index 96e91c475e..bbea0936c2 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu index 6db346c67a..8300b17446 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu index 4b1c6c89f4..a52f627d79 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu index f978a9011a..ad4a86fa37 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu index 390330ec93..62acdcb554 100644 --- a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu @@ -30,8 +30,8 @@ namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_00_generate.py b/cpp/src/neighbors/detail/cagra/search_single_cta_00_generate.py index b8f623d4c4..466afa50e4 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_00_generate.py +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_00_generate.py @@ -45,8 +45,8 @@ #define instantiate_single_cta_select_and_run( \\ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \\ template void select_and_run( \\ - raft::device_matrix_view dataset, \\ - raft::device_matrix_view graph, \\ + raft::device_matrix_view dataset, \\ + raft::device_matrix_view graph, \\ INDEX_T* const topk_indices_ptr, \\ DISTANCE_T* const topk_distances_ptr, \\ const DATA_T* const queries_ptr, \\ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu index 523f2761fc..39654d5fe6 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu index cb8b21bfe8..934c548500 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu index f5ccfa7572..fa534dbf24 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu index 1d83979a88..7f158666bf 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu index cd588e13ef..2245001657 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim128_t8.cu index b47db68273..01cbf1d313 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim128_t8.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim256_t16.cu index d875080345..2586d94f2d 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim256_t16.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim512_t32.cu index 848e71a645..082874f98a 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim512_t32.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu index de7acb56fe..741d753f7c 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu index d0e90603e2..4d609df2e4 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu index 26764c5ad9..cb345489ee 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu index 6568ab6dba..91245a4913 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu index 311f42c9a7..f065e46a9a 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu index 197aa71d7b..eb4e799420 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu index dfb47a1137..4a5ad64d14 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu index 1b874bcf9b..17fe531016 100644 --- a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu +++ b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu @@ -31,8 +31,8 @@ namespace raft::neighbors::experimental::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ diff --git a/cpp/test/neighbors/ann_cagra.cuh b/cpp/test/neighbors/ann_cagra.cuh index 9969bfd7c1..78a2d08496 100644 --- a/cpp/test/neighbors/ann_cagra.cuh +++ b/cpp/test/neighbors/ann_cagra.cuh @@ -45,7 +45,7 @@ namespace raft::neighbors::experimental::cagra { namespace { // For sort_knn_graph test template -void RandomSuffle(raft::host_matrix_view index) +void RandomSuffle(raft::host_matrix_view index) { for (IdxT i = 0; i < index.extent(0); i++) { uint64_t rand = i; @@ -65,8 +65,8 @@ void RandomSuffle(raft::host_matrix_view index) } template -testing::AssertionResult CheckOrder(raft::host_matrix_view index_test, - raft::host_matrix_view dataset) +testing::AssertionResult CheckOrder(raft::host_matrix_view index_test, + raft::host_matrix_view dataset) { for (IdxT i = 0; i < index_test.extent(0); i++) { const DatatT* const base_vec = dataset.data_handle() + i * dataset.extent(1); @@ -203,15 +203,15 @@ class AnnCagraTest : public ::testing::TestWithParam { search_params.max_queries = ps.max_queries; search_params.team_size = ps.team_size; - auto database_view = raft::make_device_matrix_view( + auto database_view = raft::make_device_matrix_view( (const DataT*)database.data(), ps.n_rows, ps.dim); { cagra::index index(handle_); if (ps.host_dataset) { - auto database_host = raft::make_host_matrix(ps.n_rows, ps.dim); + auto database_host = raft::make_host_matrix(ps.n_rows, ps.dim); raft::copy(database_host.data_handle(), database.data(), database.size(), stream_); - auto database_host_view = raft::make_host_matrix_view( + auto database_host_view = raft::make_host_matrix_view( (const DataT*)database_host.data_handle(), ps.n_rows, ps.dim); index = cagra::build(handle_, index_params, database_host_view); } else { @@ -221,12 +221,12 @@ class AnnCagraTest : public ::testing::TestWithParam { } auto index = cagra::deserialize(handle_, "cagra_index"); - auto search_queries_view = raft::make_device_matrix_view( + auto search_queries_view = raft::make_device_matrix_view( search_queries.data(), ps.n_queries, ps.dim); auto indices_out_view = - raft::make_device_matrix_view(indices_dev.data(), ps.n_queries, ps.k); - auto dists_out_view = - raft::make_device_matrix_view(distances_dev.data(), ps.n_queries, ps.k); + raft::make_device_matrix_view(indices_dev.data(), ps.n_queries, ps.k); + auto dists_out_view = raft::make_device_matrix_view( + distances_dev.data(), ps.n_queries, ps.k); cagra::search( handle_, search_params, index, search_queries_view, indices_out_view, dists_out_view); @@ -234,6 +234,7 @@ class AnnCagraTest : public ::testing::TestWithParam { update_host(indices_Cagra.data(), indices_dev.data(), queries_size, stream_); resource::sync_stream(handle_); } + // for (int i = 0; i < min(ps.n_queries, 10); i++) { // // std::cout << "query " << i << std::end; // print_vector("T", indices_naive.data() + i * ps.k, ps.k, std::cout); @@ -307,17 +308,17 @@ class AnnCagraSortTest : public ::testing::TestWithParam { { { // Step 1: Build a sorted KNN graph by CAGRA knn build - auto database_view = raft::make_device_matrix_view( + auto database_view = raft::make_device_matrix_view( (const DataT*)database.data(), ps.n_rows, ps.dim); - auto database_host = raft::make_host_matrix(ps.n_rows, ps.dim); + auto database_host = raft::make_host_matrix(ps.n_rows, ps.dim); raft::copy( database_host.data_handle(), database.data(), database.size(), handle_.get_stream()); - auto database_host_view = raft::make_host_matrix_view( + auto database_host_view = raft::make_host_matrix_view( (const DataT*)database_host.data_handle(), ps.n_rows, ps.dim); cagra::index_params index_params; auto knn_graph = - raft::make_host_matrix(ps.n_rows, index_params.intermediate_graph_degree); + raft::make_host_matrix(ps.n_rows, index_params.intermediate_graph_degree); if (ps.host_dataset) { cagra::build_knn_graph(handle_, database_host_view, knn_graph.view()); @@ -365,7 +366,6 @@ class AnnCagraSortTest : public ::testing::TestWithParam { inline std::vector generate_inputs() { - // Todo(tfeher): MULTI_CTA tests a bug, consider disabling that mode. // TODO(tfeher): test MULTI_CTA kernel with num_Parents>1 to allow multiple CTA per queries std::vector inputs = raft::util::itertools::product( {100}, diff --git a/cpp/test/neighbors/ann_cagra/search_kernel_uint64_t.cuh b/cpp/test/neighbors/ann_cagra/search_kernel_uint64_t.cuh index 562e5ac2ca..077aef5202 100644 --- a/cpp/test/neighbors/ann_cagra/search_kernel_uint64_t.cuh +++ b/cpp/test/neighbors/ann_cagra/search_kernel_uint64_t.cuh @@ -22,8 +22,8 @@ namespace raft::neighbors::experimental::cagra::detail { namespace multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ extern template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \ @@ -59,8 +59,8 @@ namespace single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ extern template void select_and_run( \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view graph, \ INDEX_T* const topk_indices_ptr, \ DISTANCE_T* const topk_distances_ptr, \ const DATA_T* const queries_ptr, \