diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 9b9475c93..303a6c6b4 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -354,6 +354,11 @@ add_library( src/spatial/knn/detail/fused_l2_knn_uint32_t_float.cu ) +target_compile_options( + cuvs INTERFACE $<$:--expt-extended-lambda + --expt-relaxed-constexpr> +) + add_library(cuvs::cuvs ALIAS cuvs) target_include_directories( diff --git a/cpp/include/cuvs/distance/detail/fused_l2_nn.cuh b/cpp/include/cuvs/distance/detail/fused_l2_nn.cuh index 4f2b251e2..f5c920e08 100644 --- a/cpp/include/cuvs/distance/detail/fused_l2_nn.cuh +++ b/cpp/include/cuvs/distance/detail/fused_l2_nn.cuh @@ -301,7 +301,7 @@ void fusedL2NNImpl(OutT* min, dim3 blk(P::Nthreads); auto nblks = raft::ceildiv(m, P::Nthreads); constexpr auto maxVal = std::numeric_limits::max(); - typedef KeyValuePair KVPair; + typedef raft::KeyValuePair KVPair; RAFT_CUDA_TRY(cudaMemsetAsync(workspace, 0, sizeof(int) * m, stream)); if (initOutBuffer) { diff --git a/cpp/include/cuvs/distance/fused_l2_nn-inl.cuh b/cpp/include/cuvs/distance/fused_l2_nn-inl.cuh index 5d170c063..c6e7acb51 100644 --- a/cpp/include/cuvs/distance/fused_l2_nn-inl.cuh +++ b/cpp/include/cuvs/distance/fused_l2_nn-inl.cuh @@ -98,33 +98,35 @@ void fusedL2NN(OutT* min, auto py = reinterpret_cast(y); if (16 % sizeof(DataT) == 0 && bytes % 16 == 0 && px % 16 == 0 && py % 16 == 0) { if (is_skinny) { - detail::fusedL2NNImpl::Policy, - ReduceOpT>( + detail::fusedL2NNImpl< + DataT, + OutT, + IdxT, + typename raft::linalg::Policy4x4Skinny::Policy, + ReduceOpT>( min, x, y, xn, yn, m, n, k, (int*)workspace, redOp, pairRedOp, sqrt, initOutBuffer, stream); } else { detail::fusedL2NNImpl::Policy, + typename raft::linalg::Policy4x4::Policy, ReduceOpT>( min, x, y, xn, yn, m, n, k, (int*)workspace, redOp, pairRedOp, sqrt, initOutBuffer, stream); } } else if (8 % sizeof(DataT) == 0 && bytes % 8 == 0 && px % 8 == 0 && py % 8 == 0) { if (is_skinny) { - detail::fusedL2NNImpl::Policy, - ReduceOpT>( + detail::fusedL2NNImpl< + DataT, + OutT, + IdxT, + typename raft::linalg::Policy4x4Skinny::Policy, + ReduceOpT>( min, x, y, xn, yn, m, n, k, (int*)workspace, redOp, pairRedOp, sqrt, initOutBuffer, stream); } else { detail::fusedL2NNImpl::Policy, + typename raft::linalg::Policy4x4::Policy, ReduceOpT>( min, x, y, xn, yn, m, n, k, (int*)workspace, redOp, pairRedOp, sqrt, initOutBuffer, stream); } @@ -133,14 +135,14 @@ void fusedL2NN(OutT* min, detail::fusedL2NNImpl::Policy, + typename raft::linalg::Policy4x4Skinny::Policy, ReduceOpT>( min, x, y, xn, yn, m, n, k, (int*)workspace, redOp, pairRedOp, sqrt, initOutBuffer, stream); } else { detail::fusedL2NNImpl::Policy, + typename raft::linalg::Policy4x4::Policy, ReduceOpT>( min, x, y, xn, yn, m, n, k, (int*)workspace, redOp, pairRedOp, sqrt, initOutBuffer, stream); } diff --git a/cpp/include/cuvs/neighbors/ball_cover-ext.cuh b/cpp/include/cuvs/neighbors/ball_cover-ext.cuh index 199444542..b1cd2b4ed 100644 --- a/cpp/include/cuvs/neighbors/ball_cover-ext.cuh +++ b/cpp/include/cuvs/neighbors/ball_cover-ext.cuh @@ -40,8 +40,8 @@ void all_knn_query(raft::resources const& handle, template void all_knn_query(raft::resources const& handle, BallCoverIndex& index, - raft::device_matrix_view inds, - raft::device_matrix_view dists, + raft::device_matrix_view inds, + raft::device_matrix_view dists, int_t k, bool perform_post_filtering = true, float weight = 1.0) RAFT_EXPLICIT; @@ -60,9 +60,9 @@ void knn_query(raft::resources const& handle, 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, + 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; @@ -91,8 +91,8 @@ void knn_query(raft::resources const& handle, 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, \ + raft::device_matrix_view inds, \ + raft::device_matrix_view dists, \ int_t k, \ bool perform_post_filtering, \ float weight); \ @@ -112,9 +112,9 @@ void knn_query(raft::resources const& handle, 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, \ + raft::device_matrix_view query, \ + raft::device_matrix_view inds, \ + raft::device_matrix_view dists, \ int_t k, \ bool perform_post_filtering, \ float weight); diff --git a/cpp/include/cuvs/neighbors/ball_cover-inl.cuh b/cpp/include/cuvs/neighbors/ball_cover-inl.cuh index bacd6a5f8..4d0f170df 100644 --- a/cpp/include/cuvs/neighbors/ball_cover-inl.cuh +++ b/cpp/include/cuvs/neighbors/ball_cover-inl.cuh @@ -204,8 +204,8 @@ void all_knn_query(raft::resources const& handle, template void all_knn_query(raft::resources const& handle, BallCoverIndex& index, - raft::device_matrix_view inds, - raft::device_matrix_view dists, + raft::device_matrix_view inds, + raft::device_matrix_view dists, int_t k, bool perform_post_filtering = true, float weight = 1.0) @@ -354,9 +354,9 @@ void knn_query(raft::resources const& handle, 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, + 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) diff --git a/cpp/include/cuvs/neighbors/ball_cover_types.hpp b/cpp/include/cuvs/neighbors/ball_cover_types.hpp index b67686fe1..c6e9fab2c 100644 --- a/cpp/include/cuvs/neighbors/ball_cover_types.hpp +++ b/cpp/include/cuvs/neighbors/ball_cover_types.hpp @@ -124,12 +124,12 @@ class BallCoverIndex { 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_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; } + raft::device_matrix_view get_X() const { return X; } cuvs::distance::DistanceType get_metric() const { return metric; } @@ -145,7 +145,7 @@ class BallCoverIndex { value_int n; value_int n_landmarks; - raft::device_matrix_view X; + raft::device_matrix_view X; cuvs::distance::DistanceType metric; @@ -158,7 +158,7 @@ class BallCoverIndex { raft::device_vector R_radius; - raft::device_matrix R; + raft::device_matrix R; protected: bool index_trained; diff --git a/cpp/include/cuvs/neighbors/brute_force-ext.cuh b/cpp/include/cuvs/neighbors/brute_force-ext.cuh index bd9a1a86e..bc4773513 100644 --- a/cpp/include/cuvs/neighbors/brute_force-ext.cuh +++ b/cpp/include/cuvs/neighbors/brute_force-ext.cuh @@ -32,25 +32,26 @@ 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, + 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, - mdspan, row_major, Accessor> dataset, - cuvs::distance::DistanceType metric = distance::DistanceType::L2Unexpanded, - T metric_arg = 0.0) RAFT_EXPLICIT; +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; + raft::device_matrix_view queries, + raft::device_matrix_view neighbors, + raft::device_matrix_view distances) RAFT_EXPLICIT; template > index, raft::device_matrix_view search, - raft::device_matrix_view indices, - raft::device_matrix_view distances, + 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, @@ -72,8 +73,8 @@ template index, raft::device_matrix_view query, - raft::device_matrix_view out_inds, - raft::device_matrix_view out_dists, + raft::device_matrix_view out_inds, + raft::device_matrix_view out_dists, cuvs::distance::DistanceType metric) RAFT_EXPLICIT; } // namespace cuvs::neighbors::brute_force @@ -89,8 +90,8 @@ void fused_l2_knn(raft::resources const& handle, raft::resources const& handle, \ std::vector> index, \ raft::device_matrix_view search, \ - raft::device_matrix_view indices, \ - raft::device_matrix_view distances, \ + raft::device_matrix_view indices, \ + raft::device_matrix_view distances, \ cuvs::distance::DistanceType metric, \ std::optional metric_arg, \ std::optional global_id_offset, \ @@ -112,32 +113,32 @@ 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); + 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); + 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, + 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, \ +#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, diff --git a/cpp/include/cuvs/neighbors/brute_force-inl.cuh b/cpp/include/cuvs/neighbors/brute_force-inl.cuh index d6eae4d4e..88904dbae 100644 --- a/cpp/include/cuvs/neighbors/brute_force-inl.cuh +++ b/cpp/include/cuvs/neighbors/brute_force-inl.cuh @@ -81,10 +81,10 @@ 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, + 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) { @@ -156,8 +156,8 @@ template > index, raft::device_matrix_view search, - raft::device_matrix_view indices, - raft::device_matrix_view distances, + 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, @@ -240,8 +240,8 @@ template index, raft::device_matrix_view query, - raft::device_matrix_view out_inds, - raft::device_matrix_view out_dists, + raft::device_matrix_view out_inds, + raft::device_matrix_view out_dists, cuvs::distance::DistanceType metric) { int k = static_cast(out_inds.extent(1)); @@ -277,7 +277,7 @@ void fused_l2_knn(raft::resources const& handle, k, rowMajorIndex, rowMajorQuery, - resource::get_cuda_stream(handle), + raft::resource::get_cuda_stream(handle), metric); } @@ -295,10 +295,11 @@ void fused_l2_knn(raft::resources const& handle, * @return the constructed brute force index */ template -index build(raft::resources const& res, - mdspan, row_major, Accessor> dataset, - cuvs::distance::DistanceType metric = distance::DistanceType::L2Unexpanded, - T metric_arg = 0.0) +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 @@ -306,7 +307,7 @@ index build(raft::resources const& res, if (metric == cuvs::distance::DistanceType::L2Expanded || metric == cuvs::distance::DistanceType::L2SqrtExpanded || metric == cuvs::distance::DistanceType::CosineExpanded) { - norms = make_device_vector(res, dataset.extent(0)); + 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, @@ -344,9 +345,9 @@ index build(raft::resources const& res, 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::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); } diff --git a/cpp/include/cuvs/neighbors/brute_force.cuh b/cpp/include/cuvs/neighbors/brute_force.cuh index 867790f02..91065d35f 100644 --- a/cpp/include/cuvs/neighbors/brute_force.cuh +++ b/cpp/include/cuvs/neighbors/brute_force.cuh @@ -83,7 +83,7 @@ template std::shared_ptr> make_batch_k_query( const raft::resources& res, const cuvs::neighbors::brute_force::index& index, - raft::device_matrix_view query, + raft::device_matrix_view query, int64_t batch_size) { return std::shared_ptr>( diff --git a/cpp/include/cuvs/neighbors/brute_force_types.hpp b/cpp/include/cuvs/neighbors/brute_force_types.hpp index 74af59ce7..6e80496f2 100644 --- a/cpp/include/cuvs/neighbors/brute_force_types.hpp +++ b/cpp/include/cuvs/neighbors/brute_force_types.hpp @@ -59,13 +59,14 @@ struct index : ann::index { /** Dataset [size, dim] */ [[nodiscard]] inline auto dataset() const noexcept - -> device_matrix_view + -> raft::device_matrix_view { return dataset_view_; } /** Dataset norms */ - [[nodiscard]] inline auto norms() const -> device_vector_view + [[nodiscard]] inline auto norms() const + -> raft::device_vector_view { return norms_view_.value(); } @@ -114,7 +115,7 @@ struct index : ann::index { * Having precomputed norms gives us a performance advantage at query time. */ index(raft::resources const& res, - raft::device_matrix_view dataset_view, + raft::device_matrix_view dataset_view, std::optional> norms_view, cuvs::distance::DistanceType metric, T metric_arg = 0.0) @@ -132,7 +133,7 @@ struct index : ann::index { * Replace the dataset with a new dataset. */ void update_dataset(raft::resources const& res, - raft::device_matrix_view dataset) + raft::device_matrix_view dataset) { dataset_view_ = dataset; } @@ -143,7 +144,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) { dataset_ = raft::make_device_matrix(dataset.extents(0), dataset.extents(1)); raft::copy(dataset_.data_handle(), @@ -154,10 +155,10 @@ struct index : ann::index { } cuvs::distance::DistanceType metric_; - raft::device_matrix dataset_; + raft::device_matrix dataset_; std::optional> norms_; std::optional> norms_view_; - raft::device_matrix_view dataset_view_; + raft::device_matrix_view dataset_view_; T metric_arg_; }; diff --git a/cpp/include/cuvs/neighbors/cagra.cuh b/cpp/include/cuvs/neighbors/cagra.cuh index f748ba595..a8e42d18a 100644 --- a/cpp/include/cuvs/neighbors/cagra.cuh +++ b/cpp/include/cuvs/neighbors/cagra.cuh @@ -74,21 +74,23 @@ namespace cuvs::neighbors::cagra { * @param[in] search_params (optional) ivf_pq search parameters */ template -void build_knn_graph(raft::resources const& res, - 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) +void build_knn_graph( + raft::resources const& res, + raft::mdspan, raft::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 = raft::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>( - dataset.data_handle(), dataset.extent(0), dataset.extent(1)); + auto dataset_internal = + raft::mdspan, raft::row_major, accessor>( + dataset.data_handle(), dataset.extent(0), dataset.extent(1)); cagra::detail::build_knn_graph( res, dataset_internal, knn_graph_internal, refine_rate, build_params, search_params); @@ -138,10 +140,11 @@ template , memory_type::device>> -void build_knn_graph(raft::resources const& res, - mdspan, row_major, accessor> dataset, - raft::host_matrix_view knn_graph, - experimental::nn_descent::index_params build_params) +void build_knn_graph( + raft::resources const& res, + raft::mdspan, raft::row_major, accessor> dataset, + raft::host_matrix_view knn_graph, + experimental::nn_descent::index_params build_params) { detail::build_knn_graph(res, dataset, knn_graph, build_params); } @@ -183,22 +186,24 @@ template , memory_type::device>, typename g_accessor = host_device_accessor, memory_type::host>> -void sort_knn_graph(raft::resources const& res, - mdspan, row_major, d_accessor> dataset, - mdspan, row_major, g_accessor> knn_graph) +void sort_knn_graph( + raft::resources const& res, + raft::mdspan, raft::row_major, d_accessor> dataset, + raft::mdspan, raft::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>( + raft::mdspan, raft::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>( - dataset.data_handle(), dataset.extent(0), dataset.extent(1)); + auto dataset_internal = + raft::mdspan, raft::row_major, d_accessor>( + dataset.data_handle(), dataset.extent(0), dataset.extent(1)); cagra::detail::graph::sort_knn_graph(res, dataset_internal, knn_graph_internal); } @@ -220,9 +225,10 @@ void sort_knn_graph(raft::resources const& res, template , memory_type::host>> -void optimize(raft::resources const& res, - mdspan, row_major, g_accessor> knn_graph, - raft::host_matrix_view new_graph) +void optimize( + raft::resources const& res, + raft::mdspan, raft::row_major, g_accessor> knn_graph, + raft::host_matrix_view new_graph) { detail::optimize(res, knn_graph, new_graph); } @@ -271,9 +277,10 @@ template , memory_type::host>> -index build(raft::resources const& res, - const index_params& params, - mdspan, row_major, Accessor> dataset) +index build( + raft::resources const& res, + const index_params& params, + raft::mdspan, raft::row_major, Accessor> dataset) { return detail::build(res, params, dataset); } @@ -299,9 +306,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), @@ -313,13 +320,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)); cagra::detail::search_main void search_with_filtering(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, CagraSampleFilterT sample_filter = CagraSampleFilterT()) { RAFT_EXPECTS( @@ -391,13 +398,13 @@ void search_with_filtering(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)); cagra::detail::search_main( diff --git a/cpp/include/cuvs/neighbors/cagra_types.hpp b/cpp/include/cuvs/neighbors/cagra_types.hpp index 987633a8a..9dee5b72a 100644 --- a/cpp/include/cuvs/neighbors/cagra_types.hpp +++ b/cpp/include/cuvs/neighbors/cagra_types.hpp @@ -259,7 +259,7 @@ struct index : ann::index { "Dataset and knn_graph must have equal number of rows"); update_dataset(res, dataset); update_graph(res, knn_graph); - resource::sync_stream(res); + raft::resource::sync_stream(res); } /** @@ -323,7 +323,7 @@ struct index : ann::index { raft::copy(graph_.data_handle(), knn_graph.data_handle(), knn_graph.size(), - resource::get_cuda_stream(res)); + raft::resource::get_cuda_stream(res)); graph_view_ = graph_.view(); } diff --git a/cpp/include/cuvs/neighbors/detail/cagra/cagra_build.cuh b/cpp/include/cuvs/neighbors/detail/cagra/cagra_build.cuh index 32db90e90..2349d5943 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/cagra_build.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/cagra_build.cuh @@ -41,12 +41,13 @@ namespace cuvs::neighbors::cagra::detail { template -void build_knn_graph(raft::resources const& res, - 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) +void build_knn_graph( + raft::resources const& res, + raft::mdspan, raft::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) { resource::detail::warn_non_pool_workspace(res, "cuvs::neighbors::cagra::build"); RAFT_EXPECTS(!build_params || build_params->metric == distance::DistanceType::L2Expanded, @@ -142,12 +143,12 @@ void build_knn_graph(raft::resources const& res, 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 + // TODO(tfeher): remove this mapping once ivf_pq accepts raft::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( + auto neighbors_view = raft::make_device_matrix_view( neighbors.data_handle(), batch.size(), neighbors.extent(1)); - auto distances_view = make_device_matrix_view( + auto distances_view = raft::make_device_matrix_view( distances.data_handle(), batch.size(), distances.extent(1)); ivf_pq::search(res, *search_params, index, queries_view, neighbors_view, distances_view); @@ -160,13 +161,13 @@ void build_knn_graph(raft::resources const& res, batch.data(), queries_view.size(), resource::get_cuda_stream(res)); - auto queries_host_view = make_host_matrix_view( + auto queries_host_view = raft::make_host_matrix_view( queries_host.data_handle(), batch.size(), batch.row_width()); - auto neighbors_host_view = make_host_matrix_view( + auto neighbors_host_view = raft::make_host_matrix_view( neighbors_host.data_handle(), batch.size(), neighbors.extent(1)); - auto refined_neighbors_host_view = make_host_matrix_view( + auto refined_neighbors_host_view = raft::make_host_matrix_view( refined_neighbors_host.data_handle(), batch.size(), top_k); - auto refined_distances_host_view = make_host_matrix_view( + auto refined_distances_host_view = raft::make_host_matrix_view( refined_distances_host.data_handle(), batch.size(), top_k); resource::sync_stream(res); @@ -178,14 +179,14 @@ void build_knn_graph(raft::resources const& res, refined_distances_host_view, build_params->metric); } else { - auto neighbor_candidates_view = make_device_matrix_view( + auto neighbor_candidates_view = raft::make_device_matrix_view( neighbors.data_handle(), batch.size(), gpu_top_k); - auto refined_neighbors_view = make_device_matrix_view( + auto refined_neighbors_view = raft::make_device_matrix_view( refined_neighbors.data_handle(), batch.size(), top_k); - auto refined_distances_view = make_device_matrix_view( + auto refined_distances_view = raft::make_device_matrix_view( refined_distances.data_handle(), batch.size(), top_k); - auto dataset_view = make_device_matrix_view( + auto dataset_view = raft::make_device_matrix_view( dataset.data_handle(), dataset.extent(0), dataset.extent(1)); cuvs::neighbors::detail::refine_device( res, @@ -242,10 +243,11 @@ void build_knn_graph(raft::resources const& res, } template -void build_knn_graph(raft::resources const& res, - mdspan, row_major, accessor> dataset, - raft::host_matrix_view knn_graph, - experimental::nn_descent::index_params build_params) +void build_knn_graph( + raft::resources const& res, + raft::mdspan, raft::row_major, accessor> dataset, + raft::host_matrix_view knn_graph, + experimental::nn_descent::index_params build_params) { auto nn_descent_idx = experimental::nn_descent::index(res, knn_graph); experimental::nn_descent::build(res, build_params, dataset, nn_descent_idx); @@ -256,7 +258,7 @@ void build_knn_graph(raft::resources const& res, host_device_accessor, g_accessor::mem_type>; auto knn_graph_internal = - mdspan, row_major, g_accessor_internal>( + raft::mdspan, raft::row_major, g_accessor_internal>( reinterpret_cast(nn_descent_idx.graph().data_handle()), nn_descent_idx.graph().extent(0), nn_descent_idx.graph().extent(1)); @@ -267,9 +269,10 @@ void build_knn_graph(raft::resources const& res, template , memory_type::host>> -void optimize(raft::resources const& res, - mdspan, row_major, g_accessor> knn_graph, - raft::host_matrix_view new_graph) +void optimize( + raft::resources const& res, + raft::mdspan, raft::row_major, g_accessor> knn_graph, + raft::host_matrix_view new_graph) { using internal_IdxT = typename std::make_unsigned::type; @@ -281,7 +284,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>( + raft::mdspan, raft::row_major, g_accessor_internal>( reinterpret_cast(knn_graph.data_handle()), knn_graph.extent(0), knn_graph.extent(1)); @@ -296,7 +299,7 @@ template build( raft::resources const& res, const index_params& params, - mdspan, row_major, Accessor> dataset, + raft::mdspan, raft::row_major, Accessor> dataset, std::optional nn_descent_params = std::nullopt, std::optional refine_rate = std::nullopt, std::optional pq_build_params = std::nullopt, diff --git a/cpp/include/cuvs/neighbors/detail/cagra/cagra_search.cuh b/cpp/include/cuvs/neighbors/detail/cagra/cagra_search.cuh index fe9da6fec..6680a8b4b 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/cagra_search.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/cagra_search.cuh @@ -105,9 +105,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, CagraSampleFilterT sample_filter = CagraSampleFilterT()) { resource::detail::warn_non_pool_workspace(res, "cuvs::neighbors::cagra::search"); @@ -150,14 +150,16 @@ void search_main(raft::resources const& res, 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)); + raft::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/cuvs/neighbors/detail/cagra/cagra_serialize.cuh b/cpp/include/cuvs/neighbors/detail/cagra/cagra_serialize.cuh index 25ab7a873..c57f03bf4 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/cagra_serialize.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/cagra_serialize.cuh @@ -69,7 +69,7 @@ void serialize(raft::resources const& res, if (include_dataset) { auto dataset = index_.dataset(); // Remove padding before saving the dataset - auto host_dataset = make_host_matrix(dataset.extent(0), dataset.extent(1)); + auto host_dataset = raft::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(), @@ -153,7 +153,7 @@ void serialize_to_hnswlib(raft::resources const& res, auto dataset = index_.dataset(); // Remove padding before saving the dataset - auto host_dataset = make_host_matrix(dataset.extent(0), dataset.extent(1)); + auto host_dataset = raft::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(), diff --git a/cpp/include/cuvs/neighbors/detail/cagra/graph_core.cuh b/cpp/include/cuvs/neighbors/detail/cagra/graph_core.cuh index b43fe0d1d..9734aa0e2 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/graph_core.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/graph_core.cuh @@ -230,9 +230,10 @@ template , memory_type::device>, typename g_accessor = host_device_accessor, memory_type::host>> -void sort_knn_graph(raft::resources const& res, - mdspan, row_major, d_accessor> dataset, - mdspan, row_major, g_accessor> knn_graph) +void sort_knn_graph( + raft::resources const& res, + raft::mdspan, raft::row_major, d_accessor> dataset, + raft::mdspan, raft::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"); @@ -317,9 +318,10 @@ void sort_knn_graph(raft::resources const& res, template , memory_type::host>> -void optimize(raft::resources const& res, - mdspan, row_major, g_accessor> knn_graph, - raft::host_matrix_view new_graph) +void optimize( + raft::resources const& res, + raft::mdspan, raft::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)); diff --git a/cpp/include/cuvs/neighbors/detail/cagra/search_multi_cta.cuh b/cpp/include/cuvs/neighbors/detail/cagra/search_multi_cta.cuh index 41a469dea..2cb11e343 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/search_multi_cta.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/search_multi_cta.cuh @@ -192,7 +192,7 @@ struct search : public search_plan_impl dataset, - raft::device_matrix_view graph, + 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/cuvs/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh b/cpp/include/cuvs/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh index 52d482ff5..27e07ae5a 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh @@ -30,7 +30,7 @@ template void select_and_run(raft::device_matrix_view dataset, - raft::device_matrix_view graph, + raft::device_matrix_view graph, INDEX_T* const topk_indices_ptr, DISTANCE_T* const topk_distances_ptr, const DATA_T* const queries_ptr, @@ -60,7 +60,7 @@ void select_and_run(raft::device_matrix_view( \ raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + 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/cuvs/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh b/cpp/include/cuvs/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh index fff67a2d4..60dc34d47 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh @@ -459,8 +459,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/cuvs/neighbors/detail/cagra/search_multi_kernel.cuh b/cpp/include/cuvs/neighbors/detail/cagra/search_multi_kernel.cuh index 86438ece0..c1d550382 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/search_multi_kernel.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/search_multi_kernel.cuh @@ -502,7 +502,7 @@ void apply_filter(INDEX_T* const result_indices_ptr, cudaStream_t cuda_stream) { const std::uint32_t block_size = 256; - const std::uint32_t grid_size = ceildiv(num_queries * result_buffer_size, block_size); + const std::uint32_t grid_size = raft::ceildiv(num_queries * result_buffer_size, block_size); apply_filter_kernel<<>>(result_indices_ptr, result_distances_ptr, @@ -674,7 +674,7 @@ 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 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/cuvs/neighbors/detail/cagra/search_plan.cuh b/cpp/include/cuvs/neighbors/detail/cagra/search_plan.cuh index 410c323d6..c4db1431a 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/search_plan.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/search_plan.cuh @@ -90,32 +90,33 @@ struct search_plan_impl : public search_plan_impl_base { int64_t graph_degree, uint32_t topk) : search_plan_impl_base(params, dim, graph_degree, topk), - hashmap(0, resource::get_cuda_stream(res)), + hashmap(0, raft::resource::get_cuda_stream(res)), num_executed_iterations(0, resource::get_cuda_stream(res)), - dev_seed(0, resource::get_cuda_stream(res)), + dev_seed(0, raft::resource::get_cuda_stream(res)), num_seeds(0) { adjust_search_params(); check_params(); calc_hashmap_params(res); set_max_dim_team(dim); - num_executed_iterations.resize(max_queries, resource::get_cuda_stream(res)); + num_executed_iterations.resize(max_queries, raft::resource::get_cuda_stream(res)); RAFT_LOG_DEBUG("# algo = %d", static_cast(algo)); } virtual ~search_plan_impl() {} - virtual void operator()(raft::resources const& res, - 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] - const std::uint32_t num_queries, - const INDEX_T* dev_seed_ptr, // [num_queries, num_seeds] - std::uint32_t* const num_executed_iterations, // [num_queries] - uint32_t topk, - SAMPLE_FILTER_T sample_filter){}; + virtual void operator()( + raft::resources const& res, + 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] + const std::uint32_t num_queries, + const INDEX_T* dev_seed_ptr, // [num_queries, num_seeds] + std::uint32_t* const num_executed_iterations, // [num_queries] + uint32_t topk, + SAMPLE_FILTER_T sample_filter){}; void adjust_search_params() { diff --git a/cpp/include/cuvs/neighbors/detail/cagra/search_single_cta.cuh b/cpp/include/cuvs/neighbors/detail/cagra/search_single_cta.cuh index c6315b292..7a2a9392c 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/search_single_cta.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/search_single_cta.cuh @@ -203,7 +203,7 @@ 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 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/cuvs/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh b/cpp/include/cuvs/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh index 38eb6d9ba..615007a9e 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh @@ -31,7 +31,7 @@ template void select_and_run( // raft::resources const& res, raft::device_matrix_view dataset, - raft::device_matrix_view graph, + 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] @@ -63,7 +63,7 @@ void select_and_run( // raft::resources const& res, extern template void \ select_and_run( \ raft::device_matrix_view dataset, \ - raft::device_matrix_view graph, \ + 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/cuvs/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh b/cpp/include/cuvs/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh index 51fabd47d..8aec44dfa 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh +++ b/cpp/include/cuvs/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh @@ -891,8 +891,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/cuvs/neighbors/detail/cagra/utils.hpp b/cpp/include/cuvs/neighbors/detail/cagra/utils.hpp index 40bab4cf7..e1cbcc878 100644 --- a/cpp/include/cuvs/neighbors/detail/cagra/utils.hpp +++ b/cpp/include/cuvs/neighbors/detail/cagra/utils.hpp @@ -155,7 +155,7 @@ struct gen_index_msb_1_mask { } // namespace utils /** - * Utility to sync memory from a host_matrix_view to a device_matrix_view + * Utility to sync memory from a host_matrix_view to a raft::device_matrix_view * * In certain situations (UVM/HMM/ATS) host memory might be directly accessible on the * device, and no extra allocations need to be performed. This class checks @@ -168,7 +168,8 @@ struct gen_index_msb_1_mask { template class device_matrix_view_from_host { public: - device_matrix_view_from_host(raft::resources const& res, host_matrix_view host_view) + device_matrix_view_from_host(raft::resources const& res, + raft::host_matrix_view host_view) : host_view_(host_view) { cudaPointerAttributes attr; @@ -181,7 +182,7 @@ class device_matrix_view_from_host { raft::copy(device_mem_->data_handle(), host_view.data_handle(), host_view.extent(0) * host_view.extent(1), - resource::get_cuda_stream(res)); + raft::resource::get_cuda_stream(res)); device_ptr = device_mem_->data_handle(); } } @@ -203,11 +204,11 @@ class device_matrix_view_from_host { }; /** - * Utility to sync memory from a device_matrix_view to a host_matrix_view + * Utility to sync memory from a raft::device_matrix_view to a host_matrix_view * * In certain situations (UVM/HMM/ATS) device memory might be directly accessible on the * host, and no extra allocations need to be performed. This class checks - * if the device_matrix_view is already accessible on the host, and only creates host + * if the raft::device_matrix_view is already accessible on the host, and only creates host * memory and copies over if necessary. In memory limited situations this is preferable * to having both a host and device copy * TODO: once the mdbuffer changes here https://github.com/wphicks/raft/blob/fea-mdbuffer @@ -230,7 +231,7 @@ class host_matrix_view_from_device { raft::copy(host_mem_->data_handle(), device_view.data_handle(), device_view.extent(0) * device_view.extent(1), - resource::get_cuda_stream(res)); + raft::resource::get_cuda_stream(res)); host_ptr = host_mem_->data_handle(); } } @@ -269,11 +270,12 @@ void copy_with_padding( raft::make_device_mdarray(res, mr, raft::make_extents(src.extent(0), padded_dim)); } if (dst.extent(1) == src.extent(1)) { - raft::copy(dst.data_handle(), src.data_handle(), src.size(), resource::get_cuda_stream(res)); + raft::copy( + dst.data_handle(), src.data_handle(), src.size(), raft::resource::get_cuda_stream(res)); } else { // copy with padding RAFT_CUDA_TRY(cudaMemsetAsync( - dst.data_handle(), 0, dst.size() * sizeof(T), resource::get_cuda_stream(res))); + dst.data_handle(), 0, dst.size() * sizeof(T), raft::resource::get_cuda_stream(res))); RAFT_CUDA_TRY(cudaMemcpy2DAsync(dst.data_handle(), sizeof(T) * dst.extent(1), src.data_handle(), @@ -281,7 +283,7 @@ void copy_with_padding( sizeof(T) * src.extent(1), src.extent(0), cudaMemcpyDefault, - resource::get_cuda_stream(res))); + raft::resource::get_cuda_stream(res))); } } } // namespace cuvs::neighbors::cagra::detail diff --git a/cpp/include/cuvs/neighbors/detail/faiss_select/MergeNetworkBlock.cuh b/cpp/include/cuvs/neighbors/detail/faiss_select/MergeNetworkBlock.cuh index 9e53c729a..14a56cfe1 100644 --- a/cpp/include/cuvs/neighbors/detail/faiss_select/MergeNetworkBlock.cuh +++ b/cpp/include/cuvs/neighbors/detail/faiss_select/MergeNetworkBlock.cuh @@ -109,7 +109,7 @@ template = WarpSize, "merge list size must be >= 32"); + static_assert(L >= raft::WarpSize, "merge list size must be >= 32"); static_assert(utils::isPowerOf2(NumThreads), "NumThreads must be a power-of-2"); static_assert(L >= NumThreads, "merge list size must be >= NumThreads"); diff --git a/cpp/include/cuvs/neighbors/detail/faiss_select/MergeNetworkWarp.cuh b/cpp/include/cuvs/neighbors/detail/faiss_select/MergeNetworkWarp.cuh index a06e282f7..a4c895d0d 100644 --- a/cpp/include/cuvs/neighbors/detail/faiss_select/MergeNetworkWarp.cuh +++ b/cpp/include/cuvs/neighbors/detail/faiss_select/MergeNetworkWarp.cuh @@ -18,15 +18,15 @@ namespace cuvs::neighbors::detail::faiss_select { // This file contains functions to: // // -perform bitonic merges on pairs of sorted lists, held in -// registers. Each list contains N * WarpSize (multiple of 32) +// registers. Each list contains N *raft::WarpSize (multiple of 32) // elements for some N. // The bitonic merge is implemented for arbitrary sizes; -// sorted list A of size N1 * WarpSize registers -// sorted list B of size N2 * WarpSize registers => -// sorted list C if size (N1 + N2) * WarpSize registers. N1 and N2 +// sorted list A of size N1 *raft::WarpSize registers +// sorted list B of size N2 *raft::WarpSize registers => +// sorted list C if size (N1 + N2) *raft::WarpSize registers. N1 and N2 // are >= 1 and don't have to be powers of 2. // -// -perform bitonic sorts on a set of N * WarpSize key/value pairs +// -perform bitonic sorts on a set of N *raft::WarpSize key/value pairs // held in registers, by using the above bitonic merge as a // primitive. // N can be an arbitrary N >= 1; i.e., the bitonic sort here supports @@ -75,7 +75,7 @@ namespace cuvs::neighbors::detail::faiss_select { // performing both < and > comparisons with the variables, so I just // stick with this. -// This function merges WarpSize / 2L lists in parallel using warp +// This function mergesraft::WarpSize / 2L lists in parallel using warp // shuffles. // It works on at most size-16 lists, as we need 32 threads for this // shuffle merge. @@ -86,7 +86,7 @@ template 15, 1 <-> 14, ... - K otherK = shfl_xor(k, 2 * L - 1); - V otherV = shfl_xor(v, 2 * L - 1); + K otherK = raft::shfl_xor(k, 2 * L - 1); + V otherV = raft::shfl_xor(v, 2 * L - 1); // Whether we are the lesser thread in the exchange bool small = !(laneId & L); @@ -117,8 +117,8 @@ inline __device__ void warpBitonicMergeLE16(K& k, V& v) #pragma unroll for (int stride = IsBitonic ? L : L / 2; stride > 0; stride /= 2) { - K otherK = shfl_xor(k, stride); - V otherV = shfl_xor(v, stride); + K otherK = raft::shfl_xor(k, stride); + V otherV = raft::shfl_xor(v, stride); // Whether we are the lesser thread in the exchange bool small = !(laneId & stride); @@ -388,8 +388,8 @@ struct BitonicMergeStep { }; /// Merges two sets of registers across the warp of any size; -/// i.e., merges a sorted k/v list of size WarpSize * N1 with a -/// sorted k/v list of size WarpSize * N2, where N1 and N2 are any +/// i.e., merges a sorted k/v list of sizeraft::WarpSize * N1 with a +/// sorted k/v list of sizeraft::WarpSize * N2, where N1 and N2 are any /// value >= 1 template inline __device__ void warpMergeAnyRegisters(K k1[N1], V v1[N1], K k2[N2], V v2[N2]) @@ -409,12 +409,12 @@ inline __device__ void warpMergeAnyRegisters(K k1[N1], V v1[N1], K k2[N2], V v2[ if (FullMerge) { // We need the other values - otherKa = shfl_xor(ka, WarpSize - 1); - otherVa = shfl_xor(va, WarpSize - 1); + otherKa = raft::shfl_xor(ka, raft::WarpSize - 1); + otherVa = raft::shfl_xor(va, raft::WarpSize - 1); } - K otherKb = shfl_xor(kb, WarpSize - 1); - V otherVb = shfl_xor(vb, WarpSize - 1); + K otherKb = raft::shfl_xor(kb, raft::WarpSize - 1); + V otherVb = raft::shfl_xor(vb, raft::WarpSize - 1); // ka is always first in the list, so we needn't use our lane // in this comparison @@ -498,8 +498,8 @@ struct BitonicSortStep { static inline __device__ void sort(K k[1], V v[1]) { // Update this code if this changes - // should go from 1 -> WarpSize in multiples of 2 - static_assert(WarpSize == 32, "unexpected warp size"); + // should go from 1 ->raft::WarpSize in multiples of 2 + static_assert(raft::WarpSize == 32, "unexpected warp size"); warpBitonicMergeLE16(k[0], v[0]); warpBitonicMergeLE16(k[0], v[0]); @@ -509,7 +509,7 @@ struct BitonicSortStep { } }; -/// Sort a list of WarpSize * N elements in registers, where N is an +/// Sort a list ofraft::WarpSize * N elements in registers, where N is an /// arbitrary >= 1 template inline __device__ void warpSortAnyRegisters(K k[N], V v[N]) diff --git a/cpp/include/cuvs/neighbors/detail/faiss_select/Select.cuh b/cpp/include/cuvs/neighbors/detail/faiss_select/Select.cuh index c92ecc77e..1ddd610ff 100644 --- a/cpp/include/cuvs/neighbors/detail/faiss_select/Select.cuh +++ b/cpp/include/cuvs/neighbors/detail/faiss_select/Select.cuh @@ -40,8 +40,8 @@ struct FinalBlockMerge<2, NumThreads, K, V, NumWarpQ, Dir, Comp> { static inline __device__ void merge(K* sharedK, V* sharedV) { // Final merge doesn't need to fully merge the second list - blockMerge(sharedK, - sharedV); + blockMerge( + sharedK, sharedV); } }; @@ -49,11 +49,17 @@ template { static inline __device__ void merge(K* sharedK, V* sharedV) { - blockMerge(sharedK, - sharedV); + blockMerge(sharedK, + sharedV); // Final merge doesn't need to fully merge the second list - blockMerge( - sharedK, sharedV); + blockMerge(sharedK, sharedV); } }; @@ -61,13 +67,19 @@ template { static inline __device__ void merge(K* sharedK, V* sharedV) { - blockMerge(sharedK, - sharedV); - blockMerge(sharedK, - sharedV); - // Final merge doesn't need to fully merge the second list - blockMerge( + blockMerge(sharedK, + sharedV); + blockMerge( sharedK, sharedV); + // Final merge doesn't need to fully merge the second list + blockMerge(sharedK, sharedV); } }; @@ -81,7 +93,7 @@ template struct BlockSelect { - static constexpr int kNumWarps = ThreadsPerBlock / WarpSize; + static constexpr int kNumWarps = ThreadsPerBlock / raft::WarpSize; static constexpr int kTotalWarpSortSize = NumWarpQ; __device__ inline BlockSelect(K initKVal, V initVVal, K* smemK, V* smemV, int k) @@ -104,13 +116,13 @@ struct BlockSelect { } int laneId = raft::laneId(); - int warpId = threadIdx.x / WarpSize; + int warpId = threadIdx.x / raft::WarpSize; warpK = sharedK + warpId * kTotalWarpSortSize; warpV = sharedV + warpId * kTotalWarpSortSize; // Fill warp queue (only the actual queue space is fine, not where // we write the per-thread queues for merging) - for (int i = laneId; i < NumWarpQ; i += WarpSize) { + for (int i = laneId; i < NumWarpQ; i += raft::WarpSize) { warpK[i] = initK; warpV[i] = initV; } @@ -178,14 +190,14 @@ struct BlockSelect { // Sort all of the per-thread queues warpSortAnyRegisters(threadK, threadV); - constexpr int kNumWarpQRegisters = NumWarpQ / WarpSize; - K warpKRegisters[kNumWarpQRegisters]; + constexpr int kNumWarpQRegisters = NumWarpQ / raft::WarpSize; + K raft::warpKRegisters[kNumWarpQRegisters]; V warpVRegisters[kNumWarpQRegisters]; #pragma unroll for (int i = 0; i < kNumWarpQRegisters; ++i) { - warpKRegisters[i] = warpK[i * WarpSize + laneId]; - warpVRegisters[i] = warpV[i * WarpSize + laneId]; + warpKRegisters[i] = warpK[i * raft::WarpSize + laneId]; + warpVRegisters[i] = warpV[i * raft::WarpSize + laneId]; } warpFence(); @@ -194,13 +206,13 @@ struct BlockSelect { // per-thread queue, merge both sorted lists together, producing // one sorted list warpMergeAnyRegisters( - warpKRegisters, warpVRegisters, threadK, threadV); + raft::warpKRegisters, warpVRegisters, threadK, threadV); // Write back out the warp queue #pragma unroll for (int i = 0; i < kNumWarpQRegisters; ++i) { - warpK[i * WarpSize + laneId] = warpKRegisters[i]; - warpV[i * WarpSize + laneId] = warpVRegisters[i]; + warpK[i * raft::WarpSize + laneId] = raft::warpKRegisters[i]; + warpV[i * raft::WarpSize + laneId] = warpVRegisters[i]; } warpFence(); @@ -264,7 +276,7 @@ struct BlockSelect { /// Specialization for k == 1 (NumWarpQ == 1) template struct BlockSelect { - static constexpr int kNumWarps = ThreadsPerBlock / WarpSize; + static constexpr int kNumWarps = ThreadsPerBlock / raft::WarpSize; __device__ inline BlockSelect(K initK, V initV, K* smemK, V* smemV, int k) : threadK(initK), threadV(initV), sharedK(smemK), sharedV(smemV) @@ -289,17 +301,17 @@ struct BlockSelect { __device__ inline void reduce() { // Reduce within the warp - KeyValuePair pair(threadK, threadV); + raft::KeyValuePair pair(threadK, threadV); if (Dir) { - pair = warpReduce(pair, max_op{}); + pair = warpReduce(pair, raft::max_op{}); } else { - pair = warpReduce(pair, min_op{}); + pair = warpReduce(pair, raft::min_op{}); } // Each warp writes out a single value int laneId = raft::laneId(); - int warpId = threadIdx.x / WarpSize; + int warpId = threadIdx.x / raft::WarpSize; if (laneId == 0) { sharedK[warpId] = pair.key; @@ -358,10 +370,14 @@ template struct WarpSelect { - static constexpr int kNumWarpQRegisters = NumWarpQ / WarpSize; + static constexpr int kNumWarpQRegisters = NumWarpQ / raft::WarpSize; __device__ inline WarpSelect(K initKVal, V initVVal, int k) - : initK(initKVal), initV(initVVal), numVals(0), warpKTop(initKVal), kLane((k - 1) % WarpSize) + : initK(initKVal), + initV(initVVal), + numVals(0), + warpKTop(initKVal), + kLane((k - 1) % raft::WarpSize) { static_assert(utils::isPowerOf2(ThreadsPerBlock), "threads must be a power-of-2"); static_assert(utils::isPowerOf2(NumWarpQ), "warp queue must be power-of-2"); @@ -425,7 +441,7 @@ struct WarpSelect { } // We have to beat at least this element - warpKTop = shfl(warpK[kNumWarpQRegisters - 1], kLane); + warpKTop = raft::shfl(warpK[kNumWarpQRegisters - 1], kLane); } /// This function handles sorting and merging together the @@ -465,7 +481,7 @@ struct WarpSelect { #pragma unroll for (int i = 0; i < kNumWarpQRegisters; ++i) { - int idx = i * WarpSize + laneId; + int idx = i * raft::WarpSize + laneId; if (idx < k) { outK[idx] = warpK[i]; @@ -503,7 +519,7 @@ struct WarpSelect { /// Specialization for k == 1 (NumWarpQ == 1) template struct WarpSelect { - static constexpr int kNumWarps = ThreadsPerBlock / WarpSize; + static constexpr int kNumWarps = ThreadsPerBlock / raft::WarpSize; __device__ inline WarpSelect(K initK, V initV, int k) : threadK(initK), threadV(initV) {} @@ -525,12 +541,12 @@ struct WarpSelect { __device__ inline void reduce() { // Reduce within the warp - KeyValuePair pair(threadK, threadV); + raft::KeyValuePair pair(threadK, threadV); if (Dir) { - pair = warpReduce(pair, max_op{}); + pair = warpReduce(pair, raft::max_op{}); } else { - pair = warpReduce(pair, min_op{}); + pair = warpReduce(pair, raft::min_op{}); } threadK = pair.key; diff --git a/cpp/include/cuvs/neighbors/detail/faiss_select/key_value_block_select.cuh b/cpp/include/cuvs/neighbors/detail/faiss_select/key_value_block_select.cuh index a9095751f..4407473a8 100644 --- a/cpp/include/cuvs/neighbors/detail/faiss_select/key_value_block_select.cuh +++ b/cpp/include/cuvs/neighbors/detail/faiss_select/key_value_block_select.cuh @@ -26,7 +26,7 @@ template struct KeyValueBlockSelect { - static constexpr int kNumWarps = ThreadsPerBlock / WarpSize; + static constexpr int kNumWarps = ThreadsPerBlock / raft::WarpSize; static constexpr int kTotalWarpSortSize = NumWarpQ; __device__ inline KeyValueBlockSelect( @@ -53,13 +53,13 @@ struct KeyValueBlockSelect { } int laneId = raft::laneId(); - int warpId = threadIdx.x / WarpSize; + int warpId = threadIdx.x / raft::WarpSize; warpK = sharedK + warpId * kTotalWarpSortSize; warpV = sharedV + warpId * kTotalWarpSortSize; // Fill warp queue (only the actual queue space is fine, not where // we write the per-thread queues for merging) - for (int i = laneId; i < NumWarpQ; i += WarpSize) { + for (int i = laneId; i < NumWarpQ; i += raft::WarpSize) { warpK[i] = initK; warpV[i].key = initVk; warpV[i].value = initVv; @@ -132,15 +132,15 @@ struct KeyValueBlockSelect { // Sort all of the per-thread queues warpSortAnyRegisters, NumThreadQ, !Dir, Comp>(threadK, threadV); - constexpr int kNumWarpQRegisters = NumWarpQ / WarpSize; - K warpKRegisters[kNumWarpQRegisters]; + constexpr int kNumWarpQRegisters = NumWarpQ / raft::WarpSize; + K raft::warpKRegisters[kNumWarpQRegisters]; KeyValuePair warpVRegisters[kNumWarpQRegisters]; #pragma unroll for (int i = 0; i < kNumWarpQRegisters; ++i) { - warpKRegisters[i] = warpK[i * WarpSize + laneId]; - warpVRegisters[i].key = warpV[i * WarpSize + laneId].key; - warpVRegisters[i].value = warpV[i * WarpSize + laneId].value; + raft::warpKRegisters[i] = warpK[i * raft::WarpSize + laneId]; + warpVRegisters[i].key = warpV[i * raft::WarpSize + laneId].key; + warpVRegisters[i].value = warpV[i * raft::WarpSize + laneId].value; } warpFence(); @@ -149,14 +149,14 @@ struct KeyValueBlockSelect { // per-thread queue, merge both sorted lists together, producing // one sorted list warpMergeAnyRegisters, kNumWarpQRegisters, NumThreadQ, !Dir, Comp, false>( - warpKRegisters, warpVRegisters, threadK, threadV); + raft::warpKRegisters, warpVRegisters, threadK, threadV); // Write back out the warp queue #pragma unroll for (int i = 0; i < kNumWarpQRegisters; ++i) { - warpK[i * WarpSize + laneId] = warpKRegisters[i]; - warpV[i * WarpSize + laneId].key = warpVRegisters[i].key; - warpV[i * WarpSize + laneId].value = warpVRegisters[i].value; + warpK[i * raft::WarpSize + laneId] = raft::warpKRegisters[i]; + warpV[i * raft::WarpSize + laneId].key = warpVRegisters[i].key; + warpV[i * raft::WarpSize + laneId].value = warpVRegisters[i].value; } warpFence(); diff --git a/cpp/include/cuvs/neighbors/detail/ivf_flat_build.cuh b/cpp/include/cuvs/neighbors/detail/ivf_flat_build.cuh index 6d9cc0de3..98d387080 100644 --- a/cpp/include/cuvs/neighbors/detail/ivf_flat_build.cuh +++ b/cpp/include/cuvs/neighbors/detail/ivf_flat_build.cuh @@ -453,10 +453,10 @@ RAFT_KERNEL unpack_interleaved_list_kernel( template void pack_list_data( raft::resources const& res, - device_matrix_view codes, + raft::device_matrix_view codes, uint32_t veclen, std::variant offset_or_indices, - device_mdspan::list_extents, row_major> list_data) + device_mdspan::list_extents, raft::row_major> list_data) { uint32_t n_rows = codes.extent(0); uint32_t dim = codes.extent(1); @@ -473,10 +473,11 @@ void pack_list_data( template void unpack_list_data( raft::resources const& res, - device_mdspan::list_extents, row_major> list_data, + device_mdspan::list_extents, raft::row_major> + list_data, uint32_t veclen, std::variant offset_or_indices, - device_matrix_view codes) + raft::device_matrix_view codes) { uint32_t n_rows = codes.extent(0); uint32_t dim = codes.extent(1); diff --git a/cpp/include/cuvs/neighbors/detail/ivf_flat_interleaved_scan-inl.cuh b/cpp/include/cuvs/neighbors/detail/ivf_flat_interleaved_scan-inl.cuh index a9999b1b1..7d6f98d52 100644 --- a/cpp/include/cuvs/neighbors/detail/ivf_flat_interleaved_scan-inl.cuh +++ b/cpp/include/cuvs/neighbors/detail/ivf_flat_interleaved_scan-inl.cuh @@ -723,7 +723,7 @@ RAFT_KERNEL __launch_bounds__(kThreadsPerBlock) // The number of interleaved groups to be processed const uint32_t num_groups = - align_warp::div(list_length + align_warp::Mask); // ceildiv by power of 2 + align_warp::div(list_length + align_warp::Mask); // raft::ceildiv by power of 2 constexpr int kUnroll = WarpSize / Veclen; constexpr uint32_t kNumWarps = kThreadsPerBlock / WarpSize; @@ -799,7 +799,7 @@ uint32_t configure_launch_x(uint32_t numQueries, uint32_t n_probes, int32_t sMem &num_blocks_per_sm, func, kThreadsPerBlock, sMemSize)); size_t min_grid_size = num_sms * num_blocks_per_sm; - size_t min_grid_x = ceildiv(min_grid_size, numQueries); + size_t min_grid_x = raft::ceildiv(min_grid_size, numQueries); return min_grid_x > n_probes ? n_probes : static_cast(min_grid_x); } diff --git a/cpp/include/cuvs/neighbors/detail/ivf_flat_serialize.cuh b/cpp/include/cuvs/neighbors/detail/ivf_flat_serialize.cuh index 478a8d420..1f181b517 100644 --- a/cpp/include/cuvs/neighbors/detail/ivf_flat_serialize.cuh +++ b/cpp/include/cuvs/neighbors/detail/ivf_flat_serialize.cuh @@ -72,7 +72,7 @@ void serialize(raft::resources const& handle, std::ostream& os, const index(index_.list_sizes().extent(0)); + auto sizes_host = raft::make_host_vector(index_.list_sizes().extent(0)); copy(sizes_host.data_handle(), index_.list_sizes().data_handle(), sizes_host.size(), diff --git a/cpp/include/cuvs/neighbors/detail/ivf_pq_build.cuh b/cpp/include/cuvs/neighbors/detail/ivf_pq_build.cuh index 08f27923e..de89ea2d9 100644 --- a/cpp/include/cuvs/neighbors/detail/ivf_pq_build.cuh +++ b/cpp/include/cuvs/neighbors/detail/ivf_pq_build.cuh @@ -185,11 +185,11 @@ void select_residuals(raft::resources const& handle, raft::matrix::gather(mapping_itr, (IdxT)dim, n_rows, row_ids, n_rows, tmp.data(), stream); raft::matrix::linewise_op(handle, - make_device_matrix_view(tmp.data(), n_rows, dim), - make_device_matrix_view(tmp.data(), n_rows, dim), + raft::make_device_matrix_view(tmp.data(), n_rows, dim), + raft::make_device_matrix_view(tmp.data(), n_rows, dim), true, raft::sub_op{}, - make_device_vector_view(center, dim)); + raft::make_device_vector_view(center, dim)); float alpha = 1.0; float beta = 0.0; @@ -222,10 +222,11 @@ void flat_compute_residuals( raft::resources const& handle, float* residuals, // [n_rows, rot_dim] IdxT n_rows, - device_matrix_view rotation_matrix, // [rot_dim, dim] - device_matrix_view centers, // [n_lists, dim_ext] - const T* dataset, // [n_rows, dim] - std::variant labels, // [n_rows] + raft::device_matrix_view + rotation_matrix, // [rot_dim, dim] + raft::device_matrix_view centers, // [n_lists, dim_ext] + const T* dataset, // [n_rows, dim] + std::variant labels, // [n_rows] rmm::mr::device_memory_resource* device_memory) { auto stream = resource::get_cuda_stream(handle); @@ -376,8 +377,8 @@ void transpose_pq_centers(const resources& handle, static_assert(extents.rank() == 3); auto extents_source = make_extents(extents.extent(0), extents.extent(2), extents.extent(1)); - auto span_source = - make_mdspan(pq_centers_source, extents_source); + auto span_source = make_mdspan( + pq_centers_source, extents_source); auto pq_centers_view = raft::make_device_vector_view( index.pq_centers().data_handle(), index.pq_centers().size()); linalg::map_offset(handle, pq_centers_view, [span_source, extents] __device__(size_t i) { @@ -567,15 +568,15 @@ void train_per_cluster(raft::resources const& handle, */ template static __device__ auto reinterpret_vectors( - device_matrix_view vectors, - device_mdspan, row_major> pq_centers) - -> device_mdspan, row_major> + raft::device_matrix_view vectors, + device_mdspan, raft::row_major> pq_centers) + -> device_mdspan, raft::row_major> { const uint32_t pq_len = pq_centers.extent(1); const uint32_t pq_dim = vectors.extent(1) / pq_len; using layout_t = typename decltype(vectors)::layout_type; using accessor_t = typename decltype(vectors)::accessor_type; - return mdspan, layout_t, accessor_t>( + return raft::mdspan, layout_t, accessor_t>( vectors.data_handle(), extent_3d{vectors.extent(0), pq_dim, pq_len}); } @@ -585,14 +586,14 @@ static __device__ auto reinterpret_vectors( * the whole byte, hence one vectors uses pq_dim bytes. */ struct unpack_codes { - device_matrix_view out_codes; + raft::device_matrix_view out_codes; /** * Create a callable to be passed to `run_on_list`. * * @param[out] out_codes the destination for the read codes. */ - __device__ inline unpack_codes(device_matrix_view out_codes) + __device__ inline unpack_codes(device_matrix_view out_codes) : out_codes{out_codes} { } @@ -606,8 +607,9 @@ struct unpack_codes { template __launch_bounds__(BlockSize) RAFT_KERNEL unpack_list_data_kernel( - device_matrix_view out_codes, - device_mdspan::list_extents, row_major> in_list_data, + raft::device_matrix_view out_codes, + device_mdspan::list_extents, raft::row_major> + in_list_data, std::variant offset_or_indices) { const uint32_t pq_dim = out_codes.extent(1); @@ -625,8 +627,9 @@ __launch_bounds__(BlockSize) RAFT_KERNEL unpack_list_data_kernel( * @param[in] stream */ inline void unpack_list_data( - device_matrix_view codes, - device_mdspan::list_extents, row_major> list_data, + raft::device_matrix_view codes, + device_mdspan::list_extents, raft::row_major> + list_data, std::variant offset_or_indices, uint32_t pq_bits, rmm::cuda_stream_view stream) @@ -655,7 +658,7 @@ inline void unpack_list_data( template void unpack_list_data(raft::resources const& res, const index& index, - device_matrix_view out_codes, + raft::device_matrix_view out_codes, uint32_t label, std::variant offset_or_indices) { @@ -696,7 +699,8 @@ struct unpack_contiguous { template __launch_bounds__(BlockSize) RAFT_KERNEL unpack_contiguous_list_data_kernel( uint8_t* out_codes, - device_mdspan::list_extents, row_major> in_list_data, + device_mdspan::list_extents, raft::row_major> + in_list_data, uint32_t n_rows, uint32_t pq_dim, std::variant offset_or_indices) @@ -708,7 +712,7 @@ __launch_bounds__(BlockSize) RAFT_KERNEL unpack_contiguous_list_data_kernel( /** * Unpack flat PQ codes from an existing list by the given offset. * - * @param[out] codes flat compressed PQ codes [n_rows, ceildiv(pq_dim * pq_bits, 8)] + * @param[out] codes flat compressed PQ codes [n_rows, raft::ceildiv(pq_dim * pq_bits, 8)] * @param[in] list_data the packed ivf::list data. * @param[in] offset_or_indices how many records in the list to skip or the exact indices. * @param[in] pq_bits codebook size (1 << pq_bits) @@ -716,7 +720,8 @@ __launch_bounds__(BlockSize) RAFT_KERNEL unpack_contiguous_list_data_kernel( */ inline void unpack_contiguous_list_data( uint8_t* codes, - device_mdspan::list_extents, row_major> list_data, + device_mdspan::list_extents, raft::row_major> + list_data, uint32_t n_rows, uint32_t pq_dim, std::variant offset_or_indices, @@ -766,9 +771,9 @@ struct reconstruct_vectors { codebook_gen codebook_kind; uint32_t cluster_ix; uint32_t pq_len; - device_mdspan, row_major> pq_centers; - device_mdspan, row_major> centers_rot; - device_mdspan, row_major> out_vectors; + device_mdspan, raft::row_major> pq_centers; + device_mdspan, raft::row_major> centers_rot; + device_mdspan, raft::row_major> out_vectors; /** * Create a callable to be passed to `run_on_list`. @@ -780,9 +785,9 @@ struct reconstruct_vectors { * @param[in] cluster_ix label/id of the cluster. */ __device__ inline reconstruct_vectors( - device_matrix_view out_vectors, - device_mdspan, row_major> pq_centers, - device_matrix_view centers_rot, + raft::device_matrix_view out_vectors, + device_mdspan, raft::row_major> pq_centers, + raft::device_matrix_view centers_rot, codebook_gen codebook_kind, uint32_t cluster_ix) : codebook_kind{codebook_kind}, @@ -818,10 +823,11 @@ struct reconstruct_vectors { template __launch_bounds__(BlockSize) RAFT_KERNEL reconstruct_list_data_kernel( - device_matrix_view out_vectors, - device_mdspan::list_extents, row_major> in_list_data, - device_mdspan, row_major> pq_centers, - device_matrix_view centers_rot, + raft::device_matrix_view out_vectors, + device_mdspan::list_extents, raft::row_major> + in_list_data, + device_mdspan, raft::row_major> pq_centers, + raft::device_matrix_view centers_rot, codebook_gen codebook_kind, uint32_t cluster_ix, std::variant offset_or_indices) @@ -837,7 +843,7 @@ __launch_bounds__(BlockSize) RAFT_KERNEL reconstruct_list_data_kernel( template void reconstruct_list_data(raft::resources const& res, const index& index, - device_matrix_view out_vectors, + raft::device_matrix_view out_vectors, uint32_t label, std::variant offset_or_indices) { @@ -852,7 +858,7 @@ void reconstruct_list_data(raft::resources const& res, "offset + output size must be not bigger than the cluster size."); } - auto tmp = make_device_mdarray( + auto tmp = raft::make_device_mdarray( res, resource::get_workspace_resource(res), make_extents(n_rows, index.rot_dim())); constexpr uint32_t kBlockSize = 256; @@ -909,7 +915,7 @@ void reconstruct_list_data(raft::resources const& res, linalg::map(res, out_vectors, utils::mapping{}, - make_device_matrix_view(out_float_ptr, n_rows, index.dim())); + raft::make_device_matrix_view(out_float_ptr, n_rows, index.dim())); } } @@ -919,14 +925,14 @@ void reconstruct_list_data(raft::resources const& res, * pq_dim bytes. */ struct pass_codes { - device_matrix_view codes; + raft::device_matrix_view codes; /** * Create a callable to be passed to `run_on_list`. * * @param[in] codes the source codes. */ - __device__ inline pass_codes(device_matrix_view codes) + __device__ inline pass_codes(device_matrix_view codes) : codes{codes} { } @@ -937,8 +943,8 @@ struct pass_codes { template __launch_bounds__(BlockSize) RAFT_KERNEL pack_list_data_kernel( - device_mdspan::list_extents, row_major> list_data, - device_matrix_view codes, + device_mdspan::list_extents, raft::row_major> list_data, + raft::device_matrix_view codes, std::variant offset_or_indices) { write_list( @@ -957,8 +963,8 @@ __launch_bounds__(BlockSize) RAFT_KERNEL pack_list_data_kernel( * @param[in] stream */ inline void pack_list_data( - device_mdspan::list_extents, row_major> list_data, - device_matrix_view codes, + device_mdspan::list_extents, raft::row_major> list_data, + raft::device_matrix_view codes, std::variant offset_or_indices, uint32_t pq_bits, rmm::cuda_stream_view stream) @@ -986,7 +992,7 @@ inline void pack_list_data( template void pack_list_data(raft::resources const& res, index* index, - device_matrix_view new_codes, + raft::device_matrix_view new_codes, uint32_t label, std::variant offset_or_indices) { @@ -1026,7 +1032,7 @@ struct pack_contiguous { template __launch_bounds__(BlockSize) RAFT_KERNEL pack_contiguous_list_data_kernel( - device_mdspan::list_extents, row_major> list_data, + device_mdspan::list_extents, raft::row_major> list_data, const uint8_t* codes, uint32_t n_rows, uint32_t pq_dim, @@ -1042,13 +1048,13 @@ __launch_bounds__(BlockSize) RAFT_KERNEL pack_contiguous_list_data_kernel( * NB: no memory allocation happens here; the list must fit the data (offset + n_rows). * * @param[out] list_data the packed ivf::list data. - * @param[in] codes flat compressed PQ codes [n_rows, ceildiv(pq_dim * pq_bits, 8)] + * @param[in] codes flat compressed PQ codes [n_rows, raft::ceildiv(pq_dim * pq_bits, 8)] * @param[in] offset_or_indices how many records in the list to skip or the exact indices. * @param[in] pq_bits codebook size (1 << pq_bits) * @param[in] stream */ inline void pack_contiguous_list_data( - device_mdspan::list_extents, row_major> list_data, + device_mdspan::list_extents, raft::row_major> list_data, const uint8_t* codes, uint32_t n_rows, uint32_t pq_dim, @@ -1118,12 +1124,12 @@ template struct encode_vectors { codebook_gen codebook_kind; uint32_t cluster_ix; - device_mdspan, row_major> pq_centers; - device_mdspan, row_major> in_vectors; + device_mdspan, raft::row_major> pq_centers; + device_mdspan, raft::row_major> in_vectors; __device__ inline encode_vectors( - device_mdspan, row_major> pq_centers, - device_matrix_view in_vectors, + device_mdspan, raft::row_major> pq_centers, + raft::device_matrix_view in_vectors, codebook_gen codebook_kind, uint32_t cluster_ix) : codebook_kind{codebook_kind}, @@ -1171,8 +1177,8 @@ struct encode_vectors { // reduce among threads #pragma unroll for (uint32_t stride = SubWarpSize >> 1; stride > 0; stride >>= 1) { - const auto other_dist = shfl_xor(min_dist, stride, SubWarpSize); - const auto other_code = shfl_xor(code, stride, SubWarpSize); + const auto other_dist = raft::shfl_xor(min_dist, stride, SubWarpSize); + const auto other_code = raft::shfl_xor(code, stride, SubWarpSize); if (other_dist < min_dist) { min_dist = other_dist; code = other_code; @@ -1184,13 +1190,13 @@ struct encode_vectors { template __launch_bounds__(BlockSize) RAFT_KERNEL process_and_fill_codes_kernel( - device_matrix_view new_vectors, + raft::device_matrix_view new_vectors, std::variant src_offset_or_indices, const uint32_t* new_labels, - device_vector_view list_sizes, - device_vector_view inds_ptrs, - device_vector_view data_ptrs, - device_mdspan, row_major> pq_centers, + raft::device_vector_view list_sizes, + raft::device_vector_view inds_ptrs, + raft::device_vector_view data_ptrs, + device_mdspan, raft::row_major> pq_centers, codebook_gen codebook_kind) { constexpr uint32_t kSubWarpSize = std::min(WarpSize, 1u << PqBits); @@ -1218,7 +1224,7 @@ __launch_bounds__(BlockSize) RAFT_KERNEL process_and_fill_codes_kernel( const uint32_t pq_dim = new_vectors.extent(1) / pq_centers.extent(1); auto pq_extents = list_spec{PqBits, pq_dim, true}.make_list_extents(out_ix + 1); auto pq_dataset = - make_mdspan(data_ptrs[cluster_ix], pq_extents); + make_mdspan(data_ptrs[cluster_ix], pq_extents); write_vector( pq_dataset, out_ix, @@ -1229,9 +1235,9 @@ __launch_bounds__(BlockSize) RAFT_KERNEL process_and_fill_codes_kernel( template __launch_bounds__(BlockSize) RAFT_KERNEL encode_list_data_kernel( - device_mdspan::list_extents, row_major> list_data, - device_matrix_view new_vectors, - device_mdspan, row_major> pq_centers, + device_mdspan::list_extents, raft::row_major> list_data, + raft::device_matrix_view new_vectors, + device_mdspan, raft::row_major> pq_centers, codebook_gen codebook_kind, uint32_t cluster_ix, std::variant offset_or_indices) @@ -1247,7 +1253,7 @@ __launch_bounds__(BlockSize) RAFT_KERNEL encode_list_data_kernel( template void encode_list_data(raft::resources const& res, index* index, - device_matrix_view new_vectors, + raft::device_matrix_view new_vectors, uint32_t label, std::variant offset_or_indices) { @@ -1257,7 +1263,7 @@ void encode_list_data(raft::resources const& res, auto mr = resource::get_workspace_resource(res); auto new_vectors_residual = - make_device_mdarray(res, mr, make_extents(n_rows, index->rot_dim())); + raft::make_device_mdarray(res, mr, make_extents(n_rows, index->rot_dim())); flat_compute_residuals(res, new_vectors_residual.data_handle(), @@ -1328,7 +1334,7 @@ void process_and_fill_codes(raft::resources const& handle, rmm::mr::device_memory_resource* mr) { auto new_vectors_residual = - make_device_mdarray(handle, mr, make_extents(n_rows, index.rot_dim())); + raft::make_device_mdarray(handle, mr, make_extents(n_rows, index.rot_dim())); flat_compute_residuals(handle, new_vectors_residual.data_handle(), @@ -1424,10 +1430,11 @@ void recompute_internal_state(const raft::resources& res, index& index) * @return offset for writing the data */ template -auto extend_list_prepare(raft::resources const& res, - index* index, - device_vector_view new_indices, - uint32_t label) -> uint32_t +auto extend_list_prepare( + raft::resources const& res, + index* index, + raft::device_vector_view new_indices, + uint32_t label) -> uint32_t { uint32_t n_rows = new_indices.extent(0); uint32_t offset; @@ -1453,11 +1460,12 @@ auto extend_list_prepare(raft::resources const& res, * See the public interface for the api and usage. */ template -void extend_list_with_codes(raft::resources const& res, - index* index, - device_matrix_view new_codes, - device_vector_view new_indices, - uint32_t label) +void extend_list_with_codes( + raft::resources const& res, + index* index, + raft::device_matrix_view new_codes, + raft::device_vector_view new_indices, + uint32_t label) { // Allocate memory and write indices auto offset = extend_list_prepare(res, index, new_indices, label); @@ -1474,8 +1482,8 @@ void extend_list_with_codes(raft::resources const& res, template void extend_list(raft::resources const& res, index* index, - device_matrix_view new_vectors, - device_vector_view new_indices, + raft::device_matrix_view new_vectors, + raft::device_vector_view new_indices, uint32_t label) { // Allocate memory and write indices diff --git a/cpp/include/cuvs/neighbors/detail/ivf_pq_codepacking.cuh b/cpp/include/cuvs/neighbors/detail/ivf_pq_codepacking.cuh index 93a3459a9..3ae826424 100644 --- a/cpp/include/cuvs/neighbors/detail/ivf_pq_codepacking.cuh +++ b/cpp/include/cuvs/neighbors/detail/ivf_pq_codepacking.cuh @@ -100,7 +100,8 @@ struct bitfield_view_t { */ template __device__ void run_on_vector( - device_mdspan::list_extents, row_major> in_list_data, + device_mdspan::list_extents, raft::row_major> + in_list_data, uint32_t in_ix, uint32_t out_ix, uint32_t pq_dim, @@ -142,7 +143,8 @@ __device__ void run_on_vector( */ template __device__ void write_vector( - device_mdspan::list_extents, row_major> out_list_data, + device_mdspan::list_extents, raft::row_major> + out_list_data, uint32_t out_ix, IdxT in_ix, uint32_t pq_dim, @@ -177,7 +179,8 @@ __device__ void write_vector( /** Process the given indices or a block of a single list (cluster). */ template __device__ void run_on_list( - device_mdspan::list_extents, row_major> in_list_data, + device_mdspan::list_extents, raft::row_major> + in_list_data, std::variant offset_or_indices, uint32_t len, uint32_t pq_dim, @@ -194,7 +197,8 @@ __device__ void run_on_list( /** Process the given indices or a block of a single list (cluster). */ template __device__ void write_list( - device_mdspan::list_extents, row_major> out_list_data, + device_mdspan::list_extents, raft::row_major> + out_list_data, std::variant offset_or_indices, uint32_t len, uint32_t pq_dim, diff --git a/cpp/include/cuvs/neighbors/detail/ivf_pq_search.cuh b/cpp/include/cuvs/neighbors/detail/ivf_pq_search.cuh index 9ef41e239..dbaf36adf 100644 --- a/cpp/include/cuvs/neighbors/detail/ivf_pq_search.cuh +++ b/cpp/include/cuvs/neighbors/detail/ivf_pq_search.cuh @@ -483,7 +483,7 @@ void ivfpq_search_worker(raft::resources const& handle, // possible. index_list_sorted_buf.resize(n_queries_probes, stream); auto index_list_buf = - make_device_mdarray(handle, mr, make_extents(n_queries_probes)); + raft::make_device_mdarray(handle, mr, make_extents(n_queries_probes)); rmm::device_uvector cluster_labels_out(n_queries_probes, stream, mr); auto index_list = index_list_buf.data_handle(); index_list_sorted = index_list_sorted_buf.data(); @@ -552,7 +552,7 @@ void ivfpq_search_worker(raft::resources const& handle, float* query_kths = nullptr; if (manage_local_topk) { query_kths_buf.emplace( - make_device_mdarray(handle, mr, make_extents(n_queries))); + raft::make_device_mdarray(handle, mr, make_extents(n_queries))); linalg::map(handle, query_kths_buf->view(), raft::const_op{dummy_block_sort_t::queue_t::kDummy}); diff --git a/cpp/include/cuvs/neighbors/detail/ivf_pq_serialize.cuh b/cpp/include/cuvs/neighbors/detail/ivf_pq_serialize.cuh index 684125d6f..79d059c46 100644 --- a/cpp/include/cuvs/neighbors/detail/ivf_pq_serialize.cuh +++ b/cpp/include/cuvs/neighbors/detail/ivf_pq_serialize.cuh @@ -73,7 +73,8 @@ void serialize(raft::resources const& handle_, std::ostream& os, const index(index.list_sizes().extents()); + auto sizes_host = + raft::make_host_mdarray(index.list_sizes().extents()); copy(sizes_host.data_handle(), index.list_sizes().data_handle(), sizes_host.size(), diff --git a/cpp/include/cuvs/neighbors/detail/knn_brute_force.cuh b/cpp/include/cuvs/neighbors/detail/knn_brute_force.cuh index 80aaab668..2b75eb048 100644 --- a/cpp/include/cuvs/neighbors/detail/knn_brute_force.cuh +++ b/cpp/include/cuvs/neighbors/detail/knn_brute_force.cuh @@ -238,12 +238,12 @@ void tiled_brute_force_knn(const raft::resources& handle, matrix::select_k( handle, - raft::make_device_matrix_view( + raft::make_device_matrix_view( temp_distances.data(), current_query_size, current_centroid_size), std::nullopt, - raft::make_device_matrix_view( + raft::make_device_matrix_view( distances + i * k, current_query_size, current_k), - raft::make_device_matrix_view( + raft::make_device_matrix_view( indices + i * k, current_query_size, current_k), select_min, true); @@ -280,13 +280,13 @@ void tiled_brute_force_knn(const raft::resources& handle, // select the actual top-k items here from the temporary output matrix::select_k( handle, - raft::make_device_matrix_view( + raft::make_device_matrix_view( temp_out_distances.data(), current_query_size, temp_out_cols), - raft::make_device_matrix_view( + raft::make_device_matrix_view( temp_out_indices.data(), current_query_size, temp_out_cols), - raft::make_device_matrix_view( + raft::make_device_matrix_view( distances + i * k, current_query_size, k), - raft::make_device_matrix_view( + raft::make_device_matrix_view( indices + i * k, current_query_size, k), select_min, true); @@ -512,9 +512,9 @@ template void brute_force_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, + raft::device_matrix_view queries, + raft::device_matrix_view neighbors, + raft::device_matrix_view distances, std::optional> query_norms = std::nullopt) { RAFT_EXPECTS(neighbors.extent(1) == distances.extent(1), "Value of k must match for outputs"); diff --git a/cpp/include/cuvs/neighbors/detail/knn_brute_force_batch_k_query.cuh b/cpp/include/cuvs/neighbors/detail/knn_brute_force_batch_k_query.cuh index 27e6c54ba..f4961ff3a 100644 --- a/cpp/include/cuvs/neighbors/detail/knn_brute_force_batch_k_query.cuh +++ b/cpp/include/cuvs/neighbors/detail/knn_brute_force_batch_k_query.cuh @@ -26,7 +26,7 @@ class gpu_batch_k_query : public batch_k_query { public: gpu_batch_k_query(const raft::resources& res, const cuvs::neighbors::brute_force::index& index, - raft::device_matrix_view query, + raft::device_matrix_view query, int64_t batch_size) : batch_k_query(res, index.size(), query.extent(0), batch_size), index(index), @@ -38,7 +38,7 @@ class gpu_batch_k_query : public batch_k_query { if (metric == cuvs::distance::DistanceType::L2Expanded || metric == cuvs::distance::DistanceType::L2SqrtExpanded || metric == cuvs::distance::DistanceType::CosineExpanded) { - query_norms = make_device_vector(res, query.extent(0)); + query_norms = raft::make_device_vector(res, query.extent(0)); if (metric == cuvs::distance::DistanceType::CosineExpanded) { raft::linalg::norm(res, @@ -92,7 +92,7 @@ class gpu_batch_k_query : public batch_k_query { } const cuvs::neighbors::brute_force::index& index; - raft::device_matrix_view query; + raft::device_matrix_view query; std::optional> query_norms; }; } // namespace cuvs::neighbors::brute_force::detail diff --git a/cpp/include/cuvs/neighbors/detail/nn_descent.cuh b/cpp/include/cuvs/neighbors/detail/nn_descent.cuh index cc7f87273..cd2208bfa 100644 --- a/cpp/include/cuvs/neighbors/detail/nn_descent.cuh +++ b/cpp/include/cuvs/neighbors/detail/nn_descent.cuh @@ -147,7 +147,7 @@ using align32 = raft::Pow2<32>; template int get_batch_size(const int it_now, const T nrow, const int batch_size) { - int it_total = ceildiv(nrow, batch_size); + int it_total = raft::ceildiv(nrow, batch_size); return (it_now == it_total - 1) ? nrow - it_now * batch_size : batch_size; } @@ -157,7 +157,7 @@ constexpr __host__ __device__ __forceinline__ int skew_dim(int ndim) { // all "4"s are for alignment if constexpr (std::is_same::value) { - ndim = ceildiv(ndim, 4) * 4; + ndim = raft::ceildiv(ndim, 4) * 4; return ndim + (ndim % 32 == 0) * 4; } } @@ -410,7 +410,7 @@ __device__ __forceinline__ void load_vec(Data_t* vec_buffer, if constexpr (std::is_same_v or std::is_same_v or std::is_same_v) { constexpr int num_load_elems_per_warp = raft::warp_size(); - for (int step = 0; step < ceildiv(padding_dims, num_load_elems_per_warp); step++) { + for (int step = 0; step < raft::ceildiv(padding_dims, num_load_elems_per_warp); step++) { int idx = step * num_load_elems_per_warp + lane_id; if (idx < load_dims) { vec_buffer[idx] = d_vec[idx]; @@ -424,7 +424,7 @@ __device__ __forceinline__ void load_vec(Data_t* vec_buffer, load_dims % 4 == 0 && padding_dims % 4 == 0) { constexpr int num_load_elems_per_warp = raft::warp_size() * 4; #pragma unroll - for (int step = 0; step < ceildiv(padding_dims, num_load_elems_per_warp); step++) { + for (int step = 0; step < raft::ceildiv(padding_dims, num_load_elems_per_warp); step++) { int idx_in_vec = step * num_load_elems_per_warp + lane_id * 4; if (idx_in_vec + 4 <= load_dims) { *(float2*)(vec_buffer + idx_in_vec) = *(float2*)(d_vec + idx_in_vec); @@ -434,7 +434,7 @@ __device__ __forceinline__ void load_vec(Data_t* vec_buffer, } } else { constexpr int num_load_elems_per_warp = raft::warp_size(); - for (int step = 0; step < ceildiv(padding_dims, num_load_elems_per_warp); step++) { + for (int step = 0; step < raft::ceildiv(padding_dims, num_load_elems_per_warp); step++) { int idx = step * num_load_elems_per_warp + lane_id; if (idx < load_dims) { vec_buffer[idx] = d_vec[idx]; @@ -464,7 +464,7 @@ RAFT_KERNEL preprocess_data_kernel(const Data_t* input_data, if (threadIdx.x == 0) { l2_norm = 0; } __syncthreads(); int lane_id = threadIdx.x % raft::warp_size(); - for (int step = 0; step < ceildiv(dim, raft::warp_size()); step++) { + for (int step = 0; step < raft::ceildiv(dim, raft::warp_size()); step++) { int idx = step * raft::warp_size() + lane_id; float part_dist = 0; if (idx < dim) { @@ -479,7 +479,7 @@ RAFT_KERNEL preprocess_data_kernel(const Data_t* input_data, __syncwarp(); } - for (int step = 0; step < ceildiv(dim, raft::warp_size()); step++) { + for (int step = 0; step < raft::ceildiv(dim, raft::warp_size()); step++) { int idx = step * raft::warp_size() + threadIdx.x; if (idx < dim) { if (l2_norms == nullptr) { @@ -528,7 +528,7 @@ __device__ void insert_to_global_graph(ResultItem elem, size_t global_idx_base = list_id * node_degree; if (elem.id() == list_id) return; - const int num_segments = ceildiv(node_degree, raft::warp_size()); + const int num_segments = raft::ceildiv(node_degree, raft::warp_size()); int loop_flag = 0; do { @@ -785,8 +785,8 @@ __launch_bounds__(BLOCK_SIZE, 4) wmma::fragment b_frag; wmma::fragment c_frag; wmma::fill_fragment(c_frag, 0.0); - for (int step = 0; step < ceildiv(data_dim, TILE_COL_WIDTH); step++) { - int num_load_elems = (step == ceildiv(data_dim, TILE_COL_WIDTH) - 1) + for (int step = 0; step < raft::ceildiv(data_dim, TILE_COL_WIDTH); step++) { + int num_load_elems = (step == raft::ceildiv(data_dim, TILE_COL_WIDTH) - 1) ? data_dim - step * TILE_COL_WIDTH : TILE_COL_WIDTH; #pragma unroll @@ -835,7 +835,7 @@ __launch_bounds__(BLOCK_SIZE, 4) } __syncthreads(); - for (int step = 0; step < ceildiv(list_new_size, num_warps); step++) { + for (int step = 0; step < raft::ceildiv(list_new_size, num_warps); step++) { int idx_in_list = step * num_warps + tx / raft::warp_size(); if (idx_in_list >= list_new_size) continue; auto min_elem = get_min_item(s_list[idx_in_list], idx_in_list, new_neighbors, s_distances); @@ -849,8 +849,8 @@ __launch_bounds__(BLOCK_SIZE, 4) __syncthreads(); wmma::fill_fragment(c_frag, 0.0); - for (int step = 0; step < ceildiv(data_dim, TILE_COL_WIDTH); step++) { - int num_load_elems = (step == ceildiv(data_dim, TILE_COL_WIDTH) - 1) + for (int step = 0; step < raft::ceildiv(data_dim, TILE_COL_WIDTH); step++) { + int num_load_elems = (step == raft::ceildiv(data_dim, TILE_COL_WIDTH) - 1) ? data_dim - step * TILE_COL_WIDTH : TILE_COL_WIDTH; if (TILE_COL_WIDTH < data_dim) { @@ -914,7 +914,7 @@ __launch_bounds__(BLOCK_SIZE, 4) } __syncthreads(); - for (int step = 0; step < ceildiv(MAX_NUM_BI_SAMPLES * 2, num_warps); step++) { + for (int step = 0; step < raft::ceildiv(MAX_NUM_BI_SAMPLES * 2, num_warps); step++) { int idx_in_list = step * num_warps + tx / raft::warp_size(); if (idx_in_list >= list_new_size && idx_in_list < MAX_NUM_BI_SAMPLES) continue; if (idx_in_list >= MAX_NUM_BI_SAMPLES + list_old_size && idx_in_list < MAX_NUM_BI_SAMPLES * 2) @@ -1227,16 +1227,17 @@ void GNND::build(Data_t* data, const Index_t nrow, Index_t* out cuvs::spatial::knn::detail::utils::batch_load_iterator vec_batches{ data, static_cast(nrow_), build_config_.dataset_dim, batch_size, stream}; for (auto const& batch : vec_batches) { - preprocess_data_kernel<<< - batch.size(), - raft::warp_size(), - sizeof(Data_t) * ceildiv(build_config_.dataset_dim, static_cast(raft::warp_size())) * - raft::warp_size(), - stream>>>(batch.data(), - d_data_.data_handle(), - build_config_.dataset_dim, - l2_norms_.data_handle(), - batch.offset()); + preprocess_data_kernel<<(raft::warp_size())) * + raft::warp_size(), + stream>>>(batch.data(), + d_data_.data_handle(), + build_config_.dataset_dim, + l2_norms_.data_handle(), + batch.offset()); } thrust::fill(thrust::device.on(stream), @@ -1370,7 +1371,7 @@ template , memory_type::host>> void build(raft::resources const& res, const index_params& params, - mdspan, row_major, Accessor> dataset, + raft::mdspan, raft::row_major, Accessor> dataset, index& idx) { RAFT_EXPECTS(dataset.extent(0) < std::numeric_limits::max() - 1, @@ -1402,7 +1403,7 @@ void build(raft::resources const& res, size_t extended_intermediate_degree = align32::roundUp( static_cast(intermediate_degree * (intermediate_degree <= 32 ? 1.0 : 1.3))); - auto int_graph = raft::make_host_matrix( + auto int_graph = raft::make_host_matrix( dataset.extent(0), static_cast(extended_graph_degree)); BuildConfig build_config{.max_dataset_size = static_cast(dataset.extent(0)), @@ -1428,9 +1429,10 @@ template , memory_type::host>> -index build(raft::resources const& res, - const index_params& params, - mdspan, row_major, Accessor> dataset) +index build( + raft::resources const& res, + const index_params& params, + raft::mdspan, raft::row_major, Accessor> dataset) { size_t intermediate_degree = params.intermediate_graph_degree; size_t graph_degree = params.graph_degree; diff --git a/cpp/include/cuvs/neighbors/detail/refine_device.cuh b/cpp/include/cuvs/neighbors/detail/refine_device.cuh index 5508ffe61..61d7a6ecc 100644 --- a/cpp/include/cuvs/neighbors/detail/refine_device.cuh +++ b/cpp/include/cuvs/neighbors/detail/refine_device.cuh @@ -37,13 +37,14 @@ namespace cuvs::neighbors::detail { * See cuvs::neighbors::refine for docs. */ template -void refine_device(raft::resources const& handle, - raft::device_matrix_view dataset, - raft::device_matrix_view queries, - raft::device_matrix_view neighbor_candidates, - raft::device_matrix_view indices, - raft::device_matrix_view distances, - distance::DistanceType metric = distance::DistanceType::L2Unexpanded) +void refine_device( + raft::resources const& handle, + raft::device_matrix_view dataset, + raft::device_matrix_view queries, + raft::device_matrix_view neighbor_candidates, + raft::device_matrix_view indices, + raft::device_matrix_view distances, + distance::DistanceType metric = distance::DistanceType::L2Unexpanded) { matrix_idx n_candidates = neighbor_candidates.extent(1); matrix_idx n_queries = queries.extent(0); diff --git a/cpp/include/cuvs/neighbors/detail/refine_host-ext.hpp b/cpp/include/cuvs/neighbors/detail/refine_host-ext.hpp index bcf24ed8d..c2dcdd91f 100644 --- a/cpp/include/cuvs/neighbors/detail/refine_host-ext.hpp +++ b/cpp/include/cuvs/neighbors/detail/refine_host-ext.hpp @@ -28,11 +28,11 @@ namespace cuvs::neighbors::detail { template [[gnu::optimize(3), gnu::optimize("tree-vectorize")]] void refine_host( - raft::host_matrix_view dataset, - raft::host_matrix_view queries, - raft::host_matrix_view neighbor_candidates, - raft::host_matrix_view indices, - raft::host_matrix_view distances, + raft::host_matrix_view dataset, + raft::host_matrix_view queries, + raft::host_matrix_view neighbor_candidates, + raft::host_matrix_view indices, + raft::host_matrix_view distances, distance::DistanceType metric = distance::DistanceType::L2Unexpanded) RAFT_EXPLICIT; } @@ -41,11 +41,11 @@ template #define instantiate_raft_neighbors_refine(IdxT, DataT, DistanceT, ExtentsT) \ extern template void cuvs::neighbors::detail::refine_host( \ - raft::host_matrix_view dataset, \ - raft::host_matrix_view queries, \ - raft::host_matrix_view neighbor_candidates, \ - raft::host_matrix_view indices, \ - raft::host_matrix_view distances, \ + raft::host_matrix_view dataset, \ + raft::host_matrix_view queries, \ + raft::host_matrix_view neighbor_candidates, \ + raft::host_matrix_view indices, \ + raft::host_matrix_view distances, \ distance::DistanceType metric); instantiate_raft_neighbors_refine(int64_t, float, float, int64_t); diff --git a/cpp/include/cuvs/neighbors/detail/refine_host-inl.hpp b/cpp/include/cuvs/neighbors/detail/refine_host-inl.hpp index e5839e8d9..ec830e58c 100644 --- a/cpp/include/cuvs/neighbors/detail/refine_host-inl.hpp +++ b/cpp/include/cuvs/neighbors/detail/refine_host-inl.hpp @@ -27,11 +27,11 @@ namespace cuvs::neighbors::detail { template [[gnu::optimize(3), gnu::optimize("tree-vectorize")]] void refine_host_impl( - raft::host_matrix_view dataset, - raft::host_matrix_view queries, - raft::host_matrix_view neighbor_candidates, - raft::host_matrix_view indices, - raft::host_matrix_view distances) + raft::host_matrix_view dataset, + raft::host_matrix_view queries, + raft::host_matrix_view neighbor_candidates, + raft::host_matrix_view indices, + raft::host_matrix_view distances) { size_t n_queries = queries.extent(0); size_t n_rows = dataset.extent(0); @@ -111,11 +111,11 @@ struct distance_comp_inner { */ template [[gnu::optimize(3), gnu::optimize("tree-vectorize")]] void refine_host( - raft::host_matrix_view dataset, - raft::host_matrix_view queries, - raft::host_matrix_view neighbor_candidates, - raft::host_matrix_view indices, - raft::host_matrix_view distances, + raft::host_matrix_view dataset, + raft::host_matrix_view queries, + raft::host_matrix_view neighbor_candidates, + raft::host_matrix_view indices, + raft::host_matrix_view distances, distance::DistanceType metric = distance::DistanceType::L2Unexpanded) { refine_check_input(dataset.extents(), diff --git a/cpp/include/cuvs/neighbors/epsilon_neighborhood.cuh b/cpp/include/cuvs/neighbors/epsilon_neighborhood.cuh index 1f04a08be..dfa300c22 100644 --- a/cpp/include/cuvs/neighbors/epsilon_neighborhood.cuh +++ b/cpp/include/cuvs/neighbors/epsilon_neighborhood.cuh @@ -99,9 +99,9 @@ void epsUnexpL2SqNeighborhood(bool* adj, */ 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_matrix_view x, + raft::device_matrix_view y, + raft::device_matrix_view adj, raft::device_vector_view vd, value_t eps) { diff --git a/cpp/include/cuvs/neighbors/ivf_flat-ext.cuh b/cpp/include/cuvs/neighbors/ivf_flat-ext.cuh index 5b9d4b3a6..3b66a589b 100644 --- a/cpp/include/cuvs/neighbors/ivf_flat-ext.cuh +++ b/cpp/include/cuvs/neighbors/ivf_flat-ext.cuh @@ -39,13 +39,13 @@ auto build(raft::resources const& handle, template auto build(raft::resources const& handle, const index_params& params, - raft::device_matrix_view dataset) + raft::device_matrix_view dataset) -> index RAFT_EXPLICIT; template void build(raft::resources const& handle, const index_params& params, - raft::device_matrix_view dataset, + raft::device_matrix_view dataset, cuvs::neighbors::ivf_flat::index& idx) RAFT_EXPLICIT; template @@ -57,7 +57,7 @@ auto extend(raft::resources const& handle, template auto extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, + raft::device_matrix_view new_vectors, std::optional> new_indices, const index& orig_index) -> index RAFT_EXPLICIT; @@ -70,7 +70,7 @@ void extend(raft::resources const& handle, template void extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, + raft::device_matrix_view new_vectors, std::optional> new_indices, index* index) RAFT_EXPLICIT; @@ -101,42 +101,42 @@ 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, + 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; + 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, \ +#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); @@ -155,7 +155,7 @@ instantiate_raft_neighbors_ivf_flat_build(uint8_t, int64_t); \ extern template auto cuvs::neighbors::ivf_flat::extend( \ raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ + raft::device_matrix_view new_vectors, \ std::optional> new_indices, \ const cuvs::neighbors::ivf_flat::index& orig_index) \ ->cuvs::neighbors::ivf_flat::index; \ @@ -169,7 +169,7 @@ instantiate_raft_neighbors_ivf_flat_build(uint8_t, int64_t); \ extern template void cuvs::neighbors::ivf_flat::extend( \ raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ + raft::device_matrix_view new_vectors, \ std::optional> new_indices, \ cuvs::neighbors::ivf_flat::index* index); @@ -179,25 +179,25 @@ 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); +#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); diff --git a/cpp/include/cuvs/neighbors/ivf_flat-inl.cuh b/cpp/include/cuvs/neighbors/ivf_flat-inl.cuh index 53bdcfe5a..d956f060c 100644 --- a/cpp/include/cuvs/neighbors/ivf_flat-inl.cuh +++ b/cpp/include/cuvs/neighbors/ivf_flat-inl.cuh @@ -109,7 +109,7 @@ auto build(raft::resources const& handle, template auto build(raft::resources const& handle, const index_params& params, - raft::device_matrix_view dataset) -> index + raft::device_matrix_view dataset) -> index { return cuvs::neighbors::ivf_flat::detail::build(handle, params, @@ -152,7 +152,7 @@ auto build(raft::resources const& handle, template void build(raft::resources const& handle, const index_params& params, - raft::device_matrix_view dataset, + raft::device_matrix_view dataset, cuvs::neighbors::ivf_flat::index& idx) { idx = cuvs::neighbors::ivf_flat::detail::build(handle, @@ -246,7 +246,7 @@ auto extend(raft::resources const& handle, */ template auto extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, + raft::device_matrix_view new_vectors, std::optional> new_indices, const index& orig_index) -> index { @@ -328,7 +328,7 @@ void extend(raft::resources const& handle, */ template void extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, + raft::device_matrix_view new_vectors, std::optional> new_indices, index* index) { @@ -521,9 +521,9 @@ 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, + raft::device_matrix_view queries, + raft::device_matrix_view neighbors, + raft::device_matrix_view distances, IvfSampleFilterT sample_filter = IvfSampleFilterT()) { RAFT_EXPECTS( @@ -584,9 +584,9 @@ 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::device_matrix_view queries, + raft::device_matrix_view neighbors, + raft::device_matrix_view distances) { search_with_filtering(handle, params, diff --git a/cpp/include/cuvs/neighbors/ivf_flat_helpers.cuh b/cpp/include/cuvs/neighbors/ivf_flat_helpers.cuh index 3baf7f8ec..25b84ec55 100644 --- a/cpp/include/cuvs/neighbors/ivf_flat_helpers.cuh +++ b/cpp/include/cuvs/neighbors/ivf_flat_helpers.cuh @@ -61,10 +61,10 @@ namespace codepacker { template void pack( raft::resources const& res, - device_matrix_view codes, + raft::device_matrix_view codes, uint32_t veclen, uint32_t offset, - device_mdspan::list_extents, row_major> list_data) + device_mdspan::list_extents, raft::row_major> list_data) { cuvs::neighbors::ivf_flat::detail::pack_list_data(res, codes, veclen, offset, list_data); } @@ -100,10 +100,11 @@ void pack( template void unpack( raft::resources const& res, - device_mdspan::list_extents, row_major> list_data, + device_mdspan::list_extents, raft::row_major> + list_data, uint32_t veclen, uint32_t offset, - device_matrix_view codes) + raft::device_matrix_view codes) { cuvs::neighbors::ivf_flat::detail::unpack_list_data( res, list_data, veclen, offset, codes); diff --git a/cpp/include/cuvs/neighbors/ivf_flat_types.hpp b/cpp/include/cuvs/neighbors/ivf_flat_types.hpp index 47b9fd9f1..e9c452f5c 100644 --- a/cpp/include/cuvs/neighbors/ivf_flat_types.hpp +++ b/cpp/include/cuvs/neighbors/ivf_flat_types.hpp @@ -89,7 +89,7 @@ static_assert(std::is_aggregate_v); template struct list_spec { using value_type = ValueT; - using list_extents = matrix_extent; + using list_extents = raft::matrix_extent; using index_type = IdxT; SizeT align_max; @@ -177,23 +177,23 @@ struct index : ann::index { * NB: This may differ from the actual list size if the shared lists have been extended by another * index */ - inline auto list_sizes() noexcept -> device_vector_view + inline auto list_sizes() noexcept -> raft::device_vector_view { return list_sizes_.view(); } [[nodiscard]] inline auto list_sizes() const noexcept - -> device_vector_view + -> raft::device_vector_view { return list_sizes_.view(); } /** k-means cluster centers corresponding to the lists [n_lists, dim] */ - inline auto centers() noexcept -> device_matrix_view + inline auto centers() noexcept -> raft::device_matrix_view { return centers_.view(); } [[nodiscard]] inline auto centers() const noexcept - -> device_matrix_view + -> raft::device_matrix_view { return centers_.view(); } @@ -274,18 +274,23 @@ struct index : ann::index { } /** Pointers to the inverted lists (clusters) data [n_lists]. */ - inline auto data_ptrs() noexcept -> device_vector_view { return data_ptrs_.view(); } - [[nodiscard]] inline auto data_ptrs() const noexcept -> device_vector_view + 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 -> device_vector_view + inline auto inds_ptrs() noexcept -> raft::device_vector_view { return inds_ptrs_.view(); } - [[nodiscard]] inline auto inds_ptrs() const noexcept -> device_vector_view + [[nodiscard]] inline auto inds_ptrs() const noexcept + -> raft::device_vector_view { return inds_ptrs_.view(); } @@ -332,7 +337,7 @@ struct index : ann::index { case cuvs::distance::DistanceType::L2SqrtExpanded: case cuvs::distance::DistanceType::L2Unexpanded: case cuvs::distance::DistanceType::L2SqrtUnexpanded: - center_norms_ = make_device_vector(res, n_lists()); + center_norms_ = raft::make_device_vector(res, n_lists()); break; default: center_norms_ = std::nullopt; } @@ -359,13 +364,13 @@ struct index : ann::index { bool adaptive_centers_; bool conservative_memory_allocation_; std::vector>> lists_; - device_vector list_sizes_; - device_matrix centers_; + raft::device_vector list_sizes_; + raft::device_matrix centers_; std::optional> center_norms_; // Computed members - device_vector data_ptrs_; - device_vector inds_ptrs_; + raft::device_vector data_ptrs_; + raft::device_vector inds_ptrs_; IdxT total_size_; /** Throw an error if the index content is inconsistent. */ diff --git a/cpp/include/cuvs/neighbors/ivf_list.hpp b/cpp/include/cuvs/neighbors/ivf_list.hpp index 8dd211266..316403893 100644 --- a/cpp/include/cuvs/neighbors/ivf_list.hpp +++ b/cpp/include/cuvs/neighbors/ivf_list.hpp @@ -51,8 +51,8 @@ list::list(raft::resources const& res, capacity = std::min(capacity, spec.align_max); } try { - data = make_device_mdarray(res, spec.make_list_extents(capacity)); - indices = make_device_vector(res, capacity); + data = raft::make_device_mdarray(res, spec.make_list_extents(capacity)); + indices = raft::make_device_vector(res, capacity); } catch (std::bad_alloc& e) { RAFT_FAIL( "ivf::list: failed to allocate a big enough list to hold all data " @@ -100,9 +100,11 @@ void resize_list(raft::resources const& res, auto new_list = std::make_shared(res, spec, new_used_size); if (old_used_size > 0) { auto copied_data_extents = spec.make_list_extents(old_used_size); - auto copied_view = - make_mdspan( - new_list->data.data_handle(), copied_data_extents); + auto copied_view = make_mdspan(new_list->data.data_handle(), copied_data_extents); copy(copied_view.data_handle(), orig_list->data.data_handle(), copied_view.size(), @@ -131,8 +133,8 @@ auto serialize_list(const raft::resources& handle, auto data_extents = store_spec.make_list_extents(size); auto data_array = - make_host_mdarray(data_extents); - auto inds_array = make_host_mdarray( + raft::make_host_mdarray(data_extents); + auto inds_array = raft::make_host_mdarray( make_extents(size)); copy(data_array.data_handle(), ld.data.data_handle(), @@ -175,8 +177,8 @@ auto deserialize_list(const raft::resources& handle, std::make_shared(handle, device_spec, size).swap(ld); auto data_extents = store_spec.make_list_extents(size); auto data_array = - make_host_mdarray(data_extents); - auto inds_array = make_host_mdarray( + raft::make_host_mdarray(data_extents); + auto inds_array = raft::make_host_mdarray( make_extents(size)); deserialize_mdspan(handle, is, data_array.view()); deserialize_mdspan(handle, is, inds_array.view()); diff --git a/cpp/include/cuvs/neighbors/ivf_list_types.hpp b/cpp/include/cuvs/neighbors/ivf_list_types.hpp index 28ecba396..67dfba807 100644 --- a/cpp/include/cuvs/neighbors/ivf_list_types.hpp +++ b/cpp/include/cuvs/neighbors/ivf_list_types.hpp @@ -46,9 +46,9 @@ struct list { using list_extents = typename spec_type::list_extents; /** Possibly encoded data; it's layout is defined by `SpecT`. */ - device_mdarray data; + device_mdarray data; /** Source indices. */ - device_mdarray, row_major> indices; + device_mdarray, raft::row_major> indices; /** The actual size of the content. */ std::atomic size; diff --git a/cpp/include/cuvs/neighbors/ivf_pq-ext.cuh b/cpp/include/cuvs/neighbors/ivf_pq-ext.cuh index d09727f2e..8c12175e5 100644 --- a/cpp/include/cuvs/neighbors/ivf_pq-ext.cuh +++ b/cpp/include/cuvs/neighbors/ivf_pq-ext.cuh @@ -31,36 +31,37 @@ namespace cuvs::neighbors::ivf_pq { template index build(raft::resources const& handle, const index_params& params, - raft::device_matrix_view dataset) RAFT_EXPLICIT; + raft::device_matrix_view dataset) RAFT_EXPLICIT; template -index extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, - std::optional> new_indices, - const index& idx) RAFT_EXPLICIT; +index extend( + raft::resources const& handle, + raft::device_matrix_view new_vectors, + std::optional> new_indices, + const index& idx) RAFT_EXPLICIT; template void extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, - std::optional> new_indices, + raft::device_matrix_view new_vectors, + std::optional> new_indices, index* idx) RAFT_EXPLICIT; template void search_with_filtering(raft::resources const& handle, 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, IvfSampleFilterT sample_filter) RAFT_EXPLICIT; template void search(raft::resources const& handle, const search_params& params, const index& idx, - raft::device_matrix_view queries, - raft::device_matrix_view neighbors, - raft::device_matrix_view distances) RAFT_EXPLICIT; + raft::device_matrix_view queries, + raft::device_matrix_view neighbors, + raft::device_matrix_view distances) RAFT_EXPLICIT; template auto build(raft::resources const& handle, @@ -139,7 +140,7 @@ search(raft::resources const& handle, extern template cuvs::neighbors::ivf_pq::index cuvs::neighbors::ivf_pq::build( \ raft::resources const& handle, \ const cuvs::neighbors::ivf_pq::index_params& params, \ - raft::device_matrix_view dataset); \ + raft::device_matrix_view dataset); \ \ extern template auto cuvs::neighbors::ivf_pq::build( \ raft::resources const& handle, \ @@ -158,14 +159,14 @@ instantiate_raft_neighbors_ivf_pq_build(uint8_t, int64_t); #define instantiate_raft_neighbors_ivf_pq_extend(T, IdxT) \ extern template cuvs::neighbors::ivf_pq::index cuvs::neighbors::ivf_pq::extend( \ raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ - std::optional> new_indices, \ + raft::device_matrix_view new_vectors, \ + std::optional> new_indices, \ const cuvs::neighbors::ivf_pq::index& idx); \ \ extern template void cuvs::neighbors::ivf_pq::extend( \ raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ - std::optional> new_indices, \ + raft::device_matrix_view new_vectors, \ + std::optional> new_indices, \ cuvs::neighbors::ivf_pq::index* idx); \ \ extern template auto cuvs::neighbors::ivf_pq::extend( \ @@ -189,34 +190,34 @@ instantiate_raft_neighbors_ivf_pq_extend(uint8_t, int64_t); #undef instantiate_raft_neighbors_ivf_pq_extend -#define instantiate_raft_neighbors_ivf_pq_search(T, IdxT) \ - extern template void cuvs::neighbors::ivf_pq::search( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::search_params& params, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - raft::device_matrix_view queries, \ - raft::device_matrix_view neighbors, \ - raft::device_matrix_view distances); \ - \ - extern template void cuvs::neighbors::ivf_pq::search( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::search_params& params, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - 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_pq::search( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::search_params& params, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - const T* queries, \ - uint32_t n_queries, \ - uint32_t k, \ - IdxT* neighbors, \ +#define instantiate_raft_neighbors_ivf_pq_search(T, IdxT) \ + extern template void cuvs::neighbors::ivf_pq::search( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::search_params& params, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + raft::device_matrix_view queries, \ + raft::device_matrix_view neighbors, \ + raft::device_matrix_view distances); \ + \ + extern template void cuvs::neighbors::ivf_pq::search( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::search_params& params, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + 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_pq::search( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::search_params& params, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + const T* queries, \ + uint32_t n_queries, \ + uint32_t k, \ + IdxT* neighbors, \ float* distances) instantiate_raft_neighbors_ivf_pq_search(float, int64_t); diff --git a/cpp/include/cuvs/neighbors/ivf_pq-inl.cuh b/cpp/include/cuvs/neighbors/ivf_pq-inl.cuh index f3c76eb86..b71a738be 100644 --- a/cpp/include/cuvs/neighbors/ivf_pq-inl.cuh +++ b/cpp/include/cuvs/neighbors/ivf_pq-inl.cuh @@ -56,7 +56,7 @@ namespace cuvs::neighbors::ivf_pq { template index build(raft::resources const& handle, const index_params& params, - raft::device_matrix_view dataset) + raft::device_matrix_view dataset) { IdxT n_rows = dataset.extent(0); IdxT dim = dataset.extent(1); @@ -78,7 +78,7 @@ index build(raft::resources const& handle, */ template index extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, + raft::device_matrix_view new_vectors, std::optional> new_indices, const index& idx) { @@ -113,7 +113,7 @@ index extend(raft::resources const& handle, */ template void extend(raft::resources const& handle, - raft::device_matrix_view new_vectors, + raft::device_matrix_view new_vectors, std::optional> new_indices, index* idx) { @@ -166,9 +166,9 @@ template void search_with_filtering(raft::resources const& handle, 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, IvfSampleFilterT sample_filter = IvfSampleFilterT{}) { RAFT_EXPECTS( @@ -222,9 +222,9 @@ template void search(raft::resources const& handle, 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) { search_with_filtering(handle, params, diff --git a/cpp/include/cuvs/neighbors/ivf_pq_helpers.cuh b/cpp/include/cuvs/neighbors/ivf_pq_helpers.cuh index 01f74a906..29c9f164b 100644 --- a/cpp/include/cuvs/neighbors/ivf_pq_helpers.cuh +++ b/cpp/include/cuvs/neighbors/ivf_pq_helpers.cuh @@ -39,7 +39,7 @@ namespace codepacker { * starting at given `offset`. * * Bit compression is removed, which means output will have pq_dim dimensional vectors (one code per - * byte, instead of ceildiv(pq_dim * pq_bits, 8) bytes of pq codes). + * byte, instead of raft::ceildiv(pq_dim * pq_bits, 8) bytes of pq codes). * * Usage example: * @code{.cpp} @@ -66,10 +66,11 @@ namespace codepacker { */ inline void unpack( raft::resources const& res, - device_mdspan::list_extents, row_major> list_data, + device_mdspan::list_extents, raft::row_major> + list_data, uint32_t pq_bits, uint32_t offset, - device_matrix_view codes) + raft::device_matrix_view codes) { ivf_pq::detail::unpack_list_data( codes, list_data, offset, pq_bits, resource::get_cuda_stream(res)); @@ -78,8 +79,8 @@ inline void unpack( /** * @brief Unpack `n_rows` consecutive records of a single list (cluster) in the compressed index * starting at given `offset`. The output codes of a single vector are contiguous, not expanded to - * one code per byte, which means the output has ceildiv(pq_dim * pq_bits, 8) bytes per PQ encoded - * vector. + * one code per byte, which means the output has raft::ceildiv(pq_dim * pq_bits, 8) bytes per PQ + * encoded vector. * * Usage example: * @code{.cpp} @@ -105,13 +106,14 @@ inline void unpack( * @param[in] n_rows How many records to unpack * @param[in] pq_dim The dimensionality of the PQ compressed records * @param[out] codes - * the destination buffer [n_rows, ceildiv(pq_dim * pq_bits, 8)]. + * the destination buffer [n_rows, raft::ceildiv(pq_dim * pq_bits, 8)]. * The length `n_rows` defines how many records to unpack, * it must be smaller than the list size. */ inline void unpack_contiguous( raft::resources const& res, - device_mdspan::list_extents, row_major> list_data, + device_mdspan::list_extents, raft::row_major> + list_data, uint32_t pq_bits, uint32_t offset, uint32_t n_rows, @@ -146,10 +148,10 @@ inline void unpack_contiguous( */ inline void pack( raft::resources const& res, - device_matrix_view codes, + raft::device_matrix_view codes, uint32_t pq_bits, uint32_t offset, - device_mdspan::list_extents, row_major> list_data) + device_mdspan::list_extents, raft::row_major> list_data) { ivf_pq::detail::pack_list_data(list_data, codes, offset, pq_bits, resource::get_cuda_stream(res)); } @@ -175,7 +177,7 @@ inline void pack( * @endcode * * @param[in] res raft resource - * @param[in] codes flat PQ codes, [n_vec, ceildiv(pq_dim * pq_bits, 8)] + * @param[in] codes flat PQ codes, [n_vec, raft::ceildiv(pq_dim * pq_bits, 8)] * @param[in] n_rows number of records * @param[in] pq_dim * @param[in] pq_bits bit length of encoded vector elements @@ -189,7 +191,7 @@ inline void pack_contiguous( uint32_t pq_dim, uint32_t pq_bits, uint32_t offset, - device_mdspan::list_extents, row_major> list_data) + device_mdspan::list_extents, raft::row_major> list_data) { ivf_pq::detail::pack_contiguous_list_data( list_data, codes, n_rows, pq_dim, offset, pq_bits, resource::get_cuda_stream(res)); @@ -223,7 +225,7 @@ inline void pack_contiguous( template void pack_list_data(raft::resources const& res, index* index, - device_matrix_view codes, + raft::device_matrix_view codes, uint32_t label, uint32_t offset) { @@ -264,7 +266,7 @@ void pack_list_data(raft::resources const& res, * * @param[in] res raft resource * @param[inout] index pointer to IVF-PQ index - * @param[in] codes flat contiguous PQ codes [n_rows, ceildiv(pq_dim * pq_bits, 8)] + * @param[in] codes flat contiguous PQ codes [n_rows, raft::ceildiv(pq_dim * pq_bits, 8)] * @param[in] n_rows how many records to pack * @param[in] label The id of the list (cluster) into which we write. * @param[in] offset how many records to skip before writing the data into the list @@ -314,7 +316,7 @@ void pack_contiguous_list_data(raft::resources const& res, template void unpack_list_data(raft::resources const& res, const index& index, - device_matrix_view out_codes, + raft::device_matrix_view out_codes, uint32_t label, uint32_t offset) { @@ -356,8 +358,8 @@ void unpack_list_data(raft::resources const& res, template void unpack_list_data(raft::resources const& res, const index& index, - device_vector_view in_cluster_indices, - device_matrix_view out_codes, + raft::device_vector_view in_cluster_indices, + raft::device_matrix_view out_codes, uint32_t label) { return ivf_pq::detail::unpack_list_data(res, index, out_codes, label, in_cluster_indices); @@ -366,7 +368,7 @@ void unpack_list_data(raft::resources const& res, /** * @brief Unpack `n_rows` consecutive PQ encoded vectors of a single list (cluster) in the * compressed index starting at given `offset`, not expanded to one code per byte. Each code in the - * output buffer occupies ceildiv(index.pq_dim() * index.pq_bits(), 8) bytes. + * output buffer occupies raft::ceildiv(index.pq_dim() * index.pq_bits(), 8) bytes. * * Usage example: * @code{.cpp} @@ -389,7 +391,7 @@ void unpack_list_data(raft::resources const& res, * @param[in] res raft resource * @param[in] index IVF-PQ index (passed by reference) * @param[out] out_codes - * the destination buffer [n_rows, ceildiv(index.pq_dim() * index.pq_bits(), 8)]. + * the destination buffer [n_rows, raft::ceildiv(index.pq_dim() * index.pq_bits(), 8)]. * The length `n_rows` defines how many records to unpack, * offset + n_rows must be smaller than or equal to the list size. * @param[in] n_rows how many codes to unpack @@ -445,7 +447,7 @@ void unpack_contiguous_list_data(raft::resources const& res, template void reconstruct_list_data(raft::resources const& res, const index& index, - device_matrix_view out_vectors, + raft::device_matrix_view out_vectors, uint32_t label, uint32_t offset) { @@ -489,8 +491,8 @@ void reconstruct_list_data(raft::resources const& res, template void reconstruct_list_data(raft::resources const& res, const index& index, - device_vector_view in_cluster_indices, - device_matrix_view out_vectors, + raft::device_vector_view in_cluster_indices, + raft::device_matrix_view out_vectors, uint32_t label) { return ivf_pq::detail::reconstruct_list_data(res, index, out_vectors, label, in_cluster_indices); @@ -509,7 +511,7 @@ void reconstruct_list_data(raft::resources const& res, * // Indices of the new vectors * auto indices = raft::make_device_vector(res, n_vec); * ... fill the indices ... - * auto new_codes = raft::make_device_matrix new_codes( + * auto new_codes = raft::make_device_matrix new_codes( * res, n_vec, index.pq_dim()); * ... fill codes ... * // extend list with new codes @@ -526,11 +528,12 @@ void reconstruct_list_data(raft::resources const& res, * @param[in] label the id of the target list (cluster). */ template -void extend_list_with_codes(raft::resources const& res, - index* index, - device_matrix_view new_codes, - device_vector_view new_indices, - uint32_t label) +void extend_list_with_codes( + raft::resources const& res, + index* index, + raft::device_matrix_view new_codes, + raft::device_vector_view new_indices, + uint32_t label) { ivf_pq::detail::extend_list_with_codes(res, index, new_codes, new_indices, label); } @@ -548,7 +551,7 @@ void extend_list_with_codes(raft::resources const& res, * // Indices of the new vectors * auto indices = raft::make_device_vector(res, n_vec); * ... fill the indices ... - * auto new_vectors = raft::make_device_matrix new_codes( + * auto new_vectors = raft::make_device_matrix new_codes( * res, n_vec, index.dim()); * ... fill vectors ... * // extend list with new vectors @@ -569,8 +572,8 @@ void extend_list_with_codes(raft::resources const& res, template void extend_list(raft::resources const& res, index* index, - device_matrix_view new_vectors, - device_vector_view new_indices, + raft::device_matrix_view new_vectors, + raft::device_vector_view new_indices, uint32_t label) { ivf_pq::detail::extend_list(res, index, new_vectors, new_indices, label); @@ -703,7 +706,7 @@ void make_rotation_matrix(raft::resources const& res, template void set_centers(raft::resources const& res, index* index, - device_matrix_view cluster_centers) + raft::device_matrix_view cluster_centers) { RAFT_EXPECTS(cluster_centers.extent(0) == index->n_lists(), "Number of rows in the new centers must be equal to the number of IVF lists"); diff --git a/cpp/include/cuvs/neighbors/ivf_pq_types.hpp b/cpp/include/cuvs/neighbors/ivf_pq_types.hpp index 0a87fd4c9..e89e11712 100644 --- a/cpp/include/cuvs/neighbors/ivf_pq_types.hpp +++ b/cpp/include/cuvs/neighbors/ivf_pq_types.hpp @@ -168,8 +168,8 @@ struct list_spec { using index_type = IdxT; /** PQ-encoded data stored in the interleaved format: * - * [ ceildiv(list_size, kIndexGroupSize) - * , ceildiv(pq_dim, (kIndexGroupVecLen * 8u) / pq_bits) + * [ raft::ceildiv(list_size, kIndexGroupSize) + * , raft::ceildiv(pq_dim, (kIndexGroupVecLen * 8u) / pq_bits) * , kIndexGroupSize * , kIndexGroupVecLen * ]. @@ -382,12 +382,12 @@ struct index : ann::index { * - codebook_gen::PER_SUBSPACE: [pq_dim , pq_len, pq_book_size] * - codebook_gen::PER_CLUSTER: [n_lists, pq_len, pq_book_size] */ - inline auto pq_centers() noexcept -> device_mdspan + inline auto pq_centers() noexcept -> device_mdspan { return pq_centers_.view(); } [[nodiscard]] inline auto pq_centers() const noexcept - -> device_mdspan + -> device_mdspan { return pq_centers_.view(); } @@ -401,36 +401,37 @@ struct index : ann::index { } /** Pointers to the inverted lists (clusters) data [n_lists]. */ - inline auto data_ptrs() noexcept -> device_vector_view + inline auto data_ptrs() noexcept -> raft::device_vector_view { return data_ptrs_.view(); } [[nodiscard]] inline auto data_ptrs() const noexcept - -> device_vector_view + -> raft::device_vector_view { - return make_mdspan( + return make_mdspan( data_ptrs_.data_handle(), data_ptrs_.extents()); } /** Pointers to the inverted lists (clusters) indices [n_lists]. */ - inline auto inds_ptrs() noexcept -> device_vector_view + inline auto inds_ptrs() noexcept -> raft::device_vector_view { return inds_ptrs_.view(); } [[nodiscard]] inline auto inds_ptrs() const noexcept - -> device_vector_view + -> raft::device_vector_view { - return make_mdspan( + return make_mdspan( inds_ptrs_.data_handle(), inds_ptrs_.extents()); } /** The transform matrix (original space -> rotated padded space) [rot_dim, dim] */ - inline auto rotation_matrix() noexcept -> device_matrix_view + inline auto rotation_matrix() noexcept + -> raft::device_matrix_view { return rotation_matrix_.view(); } [[nodiscard]] inline auto rotation_matrix() const noexcept - -> device_matrix_view + -> raft::device_matrix_view { return rotation_matrix_.view(); } @@ -444,45 +445,45 @@ struct index : ann::index { * * This span is used during search to estimate the maximum size of the workspace. */ - inline auto accum_sorted_sizes() noexcept -> host_vector_view + inline auto accum_sorted_sizes() noexcept -> host_vector_view { return accum_sorted_sizes_.view(); } [[nodiscard]] inline auto accum_sorted_sizes() const noexcept - -> host_vector_view + -> host_vector_view { return accum_sorted_sizes_.view(); } /** Sizes of the lists [n_lists]. */ - inline auto list_sizes() noexcept -> device_vector_view + inline auto list_sizes() noexcept -> raft::device_vector_view { return list_sizes_.view(); } [[nodiscard]] inline auto list_sizes() const noexcept - -> device_vector_view + -> raft::device_vector_view { return list_sizes_.view(); } /** Cluster centers corresponding to the lists in the original space [n_lists, dim_ext] */ - inline auto centers() noexcept -> device_matrix_view + inline auto centers() noexcept -> raft::device_matrix_view { return centers_.view(); } [[nodiscard]] inline auto centers() const noexcept - -> device_matrix_view + -> raft::device_matrix_view { return centers_.view(); } /** Cluster centers corresponding to the lists in the rotated space [n_lists, rot_dim] */ - inline auto centers_rot() noexcept -> device_matrix_view + inline auto centers_rot() noexcept -> raft::device_matrix_view { return centers_rot_.view(); } [[nodiscard]] inline auto centers_rot() const noexcept - -> device_matrix_view + -> raft::device_matrix_view { return centers_rot_.view(); } @@ -521,16 +522,16 @@ struct index : ann::index { // Primary data members std::vector>> lists_; - device_vector list_sizes_; - device_mdarray pq_centers_; - device_matrix centers_; - device_matrix centers_rot_; - device_matrix rotation_matrix_; + raft::device_vector list_sizes_; + device_mdarray pq_centers_; + raft::device_matrix centers_; + raft::device_matrix centers_rot_; + raft::device_matrix rotation_matrix_; // Computed members for accelerating search. - device_vector data_ptrs_; - device_vector inds_ptrs_; - host_vector accum_sorted_sizes_; + raft::device_vector data_ptrs_; + raft::device_vector inds_ptrs_; + host_vector accum_sorted_sizes_; /** Throw an error if the index content is inconsistent. */ void check_consistency() diff --git a/cpp/include/cuvs/neighbors/neighbors_types.hpp b/cpp/include/cuvs/neighbors/neighbors_types.hpp index 30da91817..c6c30655c 100644 --- a/cpp/include/cuvs/neighbors/neighbors_types.hpp +++ b/cpp/include/cuvs/neighbors/neighbors_types.hpp @@ -40,7 +40,7 @@ class batch { } /** Returns the indices for the batch */ - device_matrix_view indices() const + raft::device_matrix_view indices() const { return raft::make_const_mdspan(indices_.view()); } diff --git a/cpp/include/cuvs/neighbors/nn_descent.cuh b/cpp/include/cuvs/neighbors/nn_descent.cuh index 49ebfe0ed..0ed5cfd4a 100644 --- a/cpp/include/cuvs/neighbors/nn_descent.cuh +++ b/cpp/include/cuvs/neighbors/nn_descent.cuh @@ -58,7 +58,7 @@ namespace cuvs::neighbors::experimental::nn_descent { template index build(raft::resources const& res, index_params const& params, - raft::device_matrix_view dataset) + raft::device_matrix_view dataset) { return detail::build(res, params, dataset); } @@ -96,7 +96,7 @@ index build(raft::resources const& res, template void build(raft::resources const& res, index_params const& params, - raft::device_matrix_view dataset, + raft::device_matrix_view dataset, index& idx) { detail::build(res, params, dataset, idx); @@ -132,7 +132,7 @@ void build(raft::resources const& res, template index build(raft::resources const& res, index_params const& params, - raft::host_matrix_view dataset) + raft::host_matrix_view dataset) { return detail::build(res, params, dataset); } @@ -170,7 +170,7 @@ index build(raft::resources const& res, template void build(raft::resources const& res, index_params const& params, - raft::host_matrix_view dataset, + raft::host_matrix_view dataset, index& idx) { detail::build(res, params, dataset, idx); diff --git a/cpp/include/cuvs/neighbors/nn_descent_types.hpp b/cpp/include/cuvs/neighbors/nn_descent_types.hpp index ceb1012bd..66991755c 100644 --- a/cpp/include/cuvs/neighbors/nn_descent_types.hpp +++ b/cpp/include/cuvs/neighbors/nn_descent_types.hpp @@ -78,7 +78,7 @@ struct index : ann::index { : ann::index(), res_{res}, metric_{cuvs::distance::DistanceType::L2Expanded}, - graph_{raft::make_host_matrix(n_rows, n_cols)}, + graph_{raft::make_host_matrix(n_rows, n_cols)}, graph_view_{graph_.view()} { } @@ -98,7 +98,7 @@ struct index : ann::index { : ann::index(), res_{res}, metric_{cuvs::distance::DistanceType::L2Expanded}, - graph_{raft::make_host_matrix(0, 0)}, + graph_{raft::make_host_matrix(0, 0)}, graph_view_{graph_view} { } @@ -122,7 +122,7 @@ struct index : ann::index { } /** neighborhood graph [size, graph-degree] */ - [[nodiscard]] inline auto graph() noexcept -> host_matrix_view + [[nodiscard]] inline auto graph() noexcept -> host_matrix_view { return graph_view_; } @@ -137,8 +137,8 @@ struct index : ann::index { private: raft::resources const& res_; cuvs::distance::DistanceType metric_; - raft::host_matrix graph_; // graph to return for non-int IdxT - raft::host_matrix_view + raft::host_matrix graph_; // graph to return for non-int IdxT + raft::host_matrix_view graph_view_; // view of graph for user provided matrix }; diff --git a/cpp/include/cuvs/neighbors/refine-ext.cuh b/cpp/include/cuvs/neighbors/refine-ext.cuh index 887ead842..49bfd7301 100644 --- a/cpp/include/cuvs/neighbors/refine-ext.cuh +++ b/cpp/include/cuvs/neighbors/refine-ext.cuh @@ -30,21 +30,21 @@ namespace cuvs::neighbors { template void refine(raft::resources const& handle, - raft::device_matrix_view dataset, - raft::device_matrix_view queries, - raft::device_matrix_view neighbor_candidates, - raft::device_matrix_view indices, - raft::device_matrix_view distances, + raft::device_matrix_view dataset, + raft::device_matrix_view queries, + raft::device_matrix_view neighbor_candidates, + raft::device_matrix_view indices, + raft::device_matrix_view distances, cuvs::distance::DistanceType metric = distance::DistanceType::L2Unexpanded) RAFT_EXPLICIT; template void refine(raft::resources const& handle, - raft::host_matrix_view dataset, - raft::host_matrix_view queries, - raft::host_matrix_view neighbor_candidates, - raft::host_matrix_view indices, - raft::host_matrix_view distances, + raft::host_matrix_view dataset, + raft::host_matrix_view queries, + raft::host_matrix_view neighbor_candidates, + raft::host_matrix_view indices, + raft::host_matrix_view distances, cuvs::distance::DistanceType metric = distance::DistanceType::L2Unexpanded) RAFT_EXPLICIT; @@ -52,23 +52,23 @@ void refine(raft::resources const& handle, #endif // RAFT_EXPLICIT_INSTANTIATE_ONLY -#define instantiate_raft_neighbors_refine(idx_t, data_t, distance_t, matrix_idx) \ - extern template void cuvs::neighbors::refine( \ - raft::resources const& handle, \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view queries, \ - raft::device_matrix_view neighbor_candidates, \ - raft::device_matrix_view indices, \ - raft::device_matrix_view distances, \ - cuvs::distance::DistanceType metric); \ - \ - extern template void cuvs::neighbors::refine( \ - raft::resources const& handle, \ - raft::host_matrix_view dataset, \ - raft::host_matrix_view queries, \ - raft::host_matrix_view neighbor_candidates, \ - raft::host_matrix_view indices, \ - raft::host_matrix_view distances, \ +#define instantiate_raft_neighbors_refine(idx_t, data_t, distance_t, matrix_idx) \ + extern template void cuvs::neighbors::refine( \ + raft::resources const& handle, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view queries, \ + raft::device_matrix_view neighbor_candidates, \ + raft::device_matrix_view indices, \ + raft::device_matrix_view distances, \ + cuvs::distance::DistanceType metric); \ + \ + extern template void cuvs::neighbors::refine( \ + raft::resources const& handle, \ + raft::host_matrix_view dataset, \ + raft::host_matrix_view queries, \ + raft::host_matrix_view neighbor_candidates, \ + raft::host_matrix_view indices, \ + raft::host_matrix_view distances, \ cuvs::distance::DistanceType metric); instantiate_raft_neighbors_refine(int64_t, float, float, int64_t); diff --git a/cpp/include/cuvs/neighbors/refine-inl.cuh b/cpp/include/cuvs/neighbors/refine-inl.cuh index b42593b95..1a9b9fe34 100644 --- a/cpp/include/cuvs/neighbors/refine-inl.cuh +++ b/cpp/include/cuvs/neighbors/refine-inl.cuh @@ -68,11 +68,11 @@ namespace cuvs::neighbors { */ template void refine(raft::resources const& handle, - raft::device_matrix_view dataset, - raft::device_matrix_view queries, - raft::device_matrix_view neighbor_candidates, - raft::device_matrix_view indices, - raft::device_matrix_view distances, + raft::device_matrix_view dataset, + raft::device_matrix_view queries, + raft::device_matrix_view neighbor_candidates, + raft::device_matrix_view indices, + raft::device_matrix_view distances, distance::DistanceType metric = distance::DistanceType::L2Unexpanded) { detail::refine_device(handle, dataset, queries, neighbor_candidates, indices, distances, metric); @@ -90,11 +90,11 @@ void refine(raft::resources const& handle, */ template void refine(raft::resources const& handle, - raft::host_matrix_view dataset, - raft::host_matrix_view queries, - raft::host_matrix_view neighbor_candidates, - raft::host_matrix_view indices, - raft::host_matrix_view distances, + raft::host_matrix_view dataset, + raft::host_matrix_view queries, + raft::host_matrix_view neighbor_candidates, + raft::host_matrix_view indices, + raft::host_matrix_view distances, distance::DistanceType metric = distance::DistanceType::L2Unexpanded) { detail::refine_host(dataset, queries, neighbor_candidates, indices, distances, metric); diff --git a/cpp/include/cuvs/spatial/knn/detail/ann_utils.cuh b/cpp/include/cuvs/spatial/knn/detail/ann_utils.cuh index 8ebd24844..461479e11 100644 --- a/cpp/include/cuvs/spatial/knn/detail/ann_utils.cuh +++ b/cpp/include/cuvs/spatial/knn/detail/ann_utils.cuh @@ -285,7 +285,7 @@ void block_copy(const IdxT* in_offsets, update_host(&in_size, in_offsets + n_blocks, 1, stream); stream.synchronize(); dim3 threads(128, 1, 1); - dim3 blocks(ceildiv(in_size * n_mult, threads.x), 1, 1); + dim3 blocks(raft::ceildiv(in_size * n_mult, threads.x), 1, 1); block_copy_kernel<<>>( in_offsets, out_offsets, n_blocks, in_data, out_data, n_mult); } @@ -309,7 +309,7 @@ template void outer_add(const T* a, IdxT len_a, const T* b, IdxT len_b, T* c, rmm::cuda_stream_view stream) { dim3 threads(128, 1, 1); - dim3 blocks(ceildiv(len_a * len_b, threads.x), 1, 1); + dim3 blocks(raft::ceildiv(len_a * len_b, threads.x), 1, 1); outer_add_kernel<<>>(a, len_a, b, len_b, c); } @@ -357,7 +357,7 @@ void copy_selected(IdxT n_rows, case pointer_residency::host_and_device: case pointer_residency::device_only: { IdxT block_dim = 128; - IdxT grid_dim = ceildiv(n_rows * n_cols, block_dim); + IdxT grid_dim = raft::ceildiv(n_rows * n_cols, block_dim); copy_selected_kernel <<>>(n_rows, n_cols, src, row_ids, ld_src, dst, ld_dst); } break; diff --git a/cpp/include/cuvs/spatial/knn/detail/ball_cover.cuh b/cpp/include/cuvs/spatial/knn/detail/ball_cover.cuh index dc5f0b180..1291b160c 100644 --- a/cpp/include/cuvs/spatial/knn/detail/ball_cover.cuh +++ b/cpp/include/cuvs/spatial/knn/detail/ball_cover.cuh @@ -412,8 +412,8 @@ void rbc_all_knn_query(raft::resources const& handle, ASSERT(index.n_landmarks >= k, "number of landmark samples must be >= k"); ASSERT(!index.is_index_trained(), "index cannot be previously trained"); - rmm::device_uvector R_knn_inds(k * index.m, resource::get_cuda_stream(handle)); - rmm::device_uvector R_knn_dists(k * index.m, resource::get_cuda_stream(handle)); + rmm::device_uvector R_knn_inds(k * index.m, raft::resource::get_cuda_stream(handle)); + rmm::device_uvector R_knn_dists(k * index.m, raft::resource::get_cuda_stream(handle)); // Initialize the uvectors thrust::fill(resource::get_thrust_policy(handle), @@ -435,8 +435,9 @@ void rbc_all_knn_query(raft::resources const& handle, std::numeric_limits::max()); // For debugging / verification. Remove before releasing - rmm::device_uvector dists_counter(index.m, resource::get_cuda_stream(handle)); - rmm::device_uvector post_dists_counter(index.m, resource::get_cuda_stream(handle)); + rmm::device_uvector dists_counter(index.m, raft::resource::get_cuda_stream(handle)); + rmm::device_uvector post_dists_counter(index.m, + raft::resource::get_cuda_stream(handle)); sample_landmarks(handle, index); @@ -487,8 +488,10 @@ void rbc_knn_query(raft::resources const& handle, ASSERT(index.n_landmarks >= k, "number of landmark samples must be >= k"); ASSERT(index.is_index_trained(), "index must be previously trained"); - rmm::device_uvector R_knn_inds(k * n_query_pts, resource::get_cuda_stream(handle)); - rmm::device_uvector R_knn_dists(k * n_query_pts, resource::get_cuda_stream(handle)); + rmm::device_uvector R_knn_inds(k * n_query_pts, + raft::resource::get_cuda_stream(handle)); + rmm::device_uvector R_knn_dists(k * n_query_pts, + raft::resource::get_cuda_stream(handle)); // Initialize the uvectors thrust::fill(resource::get_thrust_policy(handle), @@ -512,8 +515,9 @@ void rbc_knn_query(raft::resources const& handle, k_closest_landmarks(handle, index, query, n_query_pts, k, R_knn_inds.data(), R_knn_dists.data()); // For debugging / verification. Remove before releasing - rmm::device_uvector dists_counter(index.m, resource::get_cuda_stream(handle)); - rmm::device_uvector post_dists_counter(index.m, resource::get_cuda_stream(handle)); + rmm::device_uvector dists_counter(index.m, raft::resource::get_cuda_stream(handle)); + rmm::device_uvector post_dists_counter(index.m, + raft::resource::get_cuda_stream(handle)); thrust::fill(resource::get_thrust_policy(handle), post_dists_counter.data(), post_dists_counter.data() + post_dists_counter.size(), diff --git a/cpp/include/cuvs/spatial/knn/detail/haversine_distance.cuh b/cpp/include/cuvs/spatial/knn/detail/haversine_distance.cuh index 7f643c62a..44fa210b8 100644 --- a/cpp/include/cuvs/spatial/knn/detail/haversine_distance.cuh +++ b/cpp/include/cuvs/spatial/knn/detail/haversine_distance.cuh @@ -60,7 +60,7 @@ RAFT_KERNEL haversine_knn_kernel(value_idx* out_inds, size_t n_index_rows, int k) { - constexpr int kNumWarps = tpb / WarpSize; + constexpr int kNumWarps = tpb / raft::WarpSize; __shared__ value_t smemK[kNumWarps * warp_q]; __shared__ value_idx smemV[kNumWarps * warp_q]; @@ -70,7 +70,7 @@ RAFT_KERNEL haversine_knn_kernel(value_idx* out_inds, std::numeric_limits::max(), std::numeric_limits::max(), smemK, smemV, k); // Grid is exactly sized to rows available - int limit = Pow2::roundDown(n_index_rows); + int limit = Pow2::roundDown(n_index_rows); const value_t* query_ptr = query + (blockIdx.x * 2); value_t x1 = query_ptr[0]; diff --git a/cpp/src/neighbors/ball_cover.cu b/cpp/src/neighbors/ball_cover.cu index f659dadfb..c9a1e9763 100644 --- a/cpp/src/neighbors/ball_cover.cu +++ b/cpp/src/neighbors/ball_cover.cu @@ -34,8 +34,8 @@ 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, \ + raft::device_matrix_view inds, \ + raft::device_matrix_view dists, \ int_t k, \ bool perform_post_filtering, \ float weight); \ @@ -54,9 +54,9 @@ 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, \ + raft::device_matrix_view query, \ + raft::device_matrix_view inds, \ + raft::device_matrix_view dists, \ int_t k, \ bool perform_post_filtering, \ float weight); diff --git a/cpp/src/neighbors/brute_force_00_generate.py b/cpp/src/neighbors/brute_force_00_generate.py index 26a56b1de..21f984502 100644 --- a/cpp/src/neighbors/brute_force_00_generate.py +++ b/cpp/src/neighbors/brute_force_00_generate.py @@ -49,8 +49,8 @@ raft::resources const& handle, \\ std::vector> index, \\ raft::device_matrix_view search, \\ - raft::device_matrix_view indices, \\ - raft::device_matrix_view distances, \\ + raft::device_matrix_view indices, \\ + raft::device_matrix_view distances, \\ cuvs::distance::DistanceType metric, \\ std::optional metric_arg, \\ std::optional global_id_offset, \\ @@ -64,8 +64,8 @@ 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, \\ + raft::device_matrix_view out_inds, \\ + raft::device_matrix_view out_dists, \\ cuvs::distance::DistanceType metric); """ diff --git a/cpp/src/neighbors/brute_force_fused_l2_knn_float_int64_t.cu b/cpp/src/neighbors/brute_force_fused_l2_knn_float_int64_t.cu index 14be76bbd..5e61f59df 100644 --- a/cpp/src/neighbors/brute_force_fused_l2_knn_float_int64_t.cu +++ b/cpp/src/neighbors/brute_force_fused_l2_knn_float_int64_t.cu @@ -27,14 +27,14 @@ #include #include -#define instantiate_raft_neighbors_brute_force_fused_l2_knn( \ - value_t, idx_t, idx_layout, query_layout) \ - 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, \ +#define instantiate_raft_neighbors_brute_force_fused_l2_knn( \ + value_t, idx_t, idx_layout, query_layout) \ + 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, diff --git a/cpp/src/neighbors/brute_force_knn_index_float.cu b/cpp/src/neighbors/brute_force_knn_index_float.cu index 5981574e9..1b98a37d0 100644 --- a/cpp/src/neighbors/brute_force_knn_index_float.cu +++ b/cpp/src/neighbors/brute_force_knn_index_float.cu @@ -21,19 +21,19 @@ template void cuvs::neighbors::brute_force::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); + raft::device_matrix_view queries, + raft::device_matrix_view neighbors, + raft::device_matrix_view distances); template void cuvs::neighbors::brute_force::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); + raft::device_matrix_view queries, + raft::device_matrix_view neighbors, + raft::device_matrix_view distances); template cuvs::neighbors::brute_force::index cuvs::neighbors::brute_force::build( raft::resources const& res, - raft::device_matrix_view dataset, + raft::device_matrix_view dataset, cuvs::distance::DistanceType metric, float metric_arg); diff --git a/cpp/src/neighbors/brute_force_knn_int64_t_float_int64_t.cu b/cpp/src/neighbors/brute_force_knn_int64_t_float_int64_t.cu index 7e30acedc..fbcabc642 100644 --- a/cpp/src/neighbors/brute_force_knn_int64_t_float_int64_t.cu +++ b/cpp/src/neighbors/brute_force_knn_int64_t_float_int64_t.cu @@ -34,8 +34,8 @@ raft::resources const& handle, \ std::vector> index, \ raft::device_matrix_view search, \ - raft::device_matrix_view indices, \ - raft::device_matrix_view distances, \ + raft::device_matrix_view indices, \ + raft::device_matrix_view distances, \ cuvs::distance::DistanceType metric, \ std::optional metric_arg, \ std::optional global_id_offset, \ diff --git a/cpp/src/neighbors/brute_force_knn_int64_t_float_uint32_t.cu b/cpp/src/neighbors/brute_force_knn_int64_t_float_uint32_t.cu index be3281f59..f9ddf2ddd 100644 --- a/cpp/src/neighbors/brute_force_knn_int64_t_float_uint32_t.cu +++ b/cpp/src/neighbors/brute_force_knn_int64_t_float_uint32_t.cu @@ -34,8 +34,8 @@ raft::resources const& handle, \ std::vector> index, \ raft::device_matrix_view search, \ - raft::device_matrix_view indices, \ - raft::device_matrix_view distances, \ + raft::device_matrix_view indices, \ + raft::device_matrix_view distances, \ cuvs::distance::DistanceType metric, \ std::optional metric_arg, \ std::optional global_id_offset, \ diff --git a/cpp/src/neighbors/brute_force_knn_int_float_int.cu b/cpp/src/neighbors/brute_force_knn_int_float_int.cu index ba94586a8..bf24dd9df 100644 --- a/cpp/src/neighbors/brute_force_knn_int_float_int.cu +++ b/cpp/src/neighbors/brute_force_knn_int_float_int.cu @@ -34,8 +34,8 @@ raft::resources const& handle, \ std::vector> index, \ raft::device_matrix_view search, \ - raft::device_matrix_view indices, \ - raft::device_matrix_view distances, \ + raft::device_matrix_view indices, \ + raft::device_matrix_view distances, \ cuvs::distance::DistanceType metric, \ std::optional metric_arg, \ std::optional global_id_offset, \ diff --git a/cpp/src/neighbors/brute_force_knn_uint32_t_float_uint32_t.cu b/cpp/src/neighbors/brute_force_knn_uint32_t_float_uint32_t.cu index 29b352613..2514352a0 100644 --- a/cpp/src/neighbors/brute_force_knn_uint32_t_float_uint32_t.cu +++ b/cpp/src/neighbors/brute_force_knn_uint32_t_float_uint32_t.cu @@ -34,8 +34,8 @@ raft::resources const& handle, \ std::vector> index, \ raft::device_matrix_view search, \ - raft::device_matrix_view indices, \ - raft::device_matrix_view distances, \ + raft::device_matrix_view indices, \ + raft::device_matrix_view distances, \ cuvs::distance::DistanceType metric, \ std::optional metric_arg, \ std::optional global_id_offset, \ 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 448db5d93..8826439c4 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 @@ -47,8 +47,8 @@ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu index eb725dc6d..c457e87c1 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 4fb908832..ab5e2821d 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 1ef7848e0..bb517f6bc 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 559f0388d..f99ac9340 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 4a30fdfce..766f65c08 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 5b733576c..d9d5ee97b 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 c7deda1b6..062817e63 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 0e1b99dfb..b6c0cd07b 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 40449a2fa..d71f0bfbc 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 cc7666ff3..041e162f6 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 fd5754e75..6e38154b5 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 073abe6c3..7663ea28d 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 f2a92dd34..dbbf3438e 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 0a9e48165..f842d6eef 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 d1ba1c3a8..ca0770533 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 a83ae6444..4d8376946 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::multi_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 fac65cc93..524500055 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 @@ -47,8 +47,8 @@ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 97733607c..0534c84fc 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 a537c3ba2..c29a1ca66 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 9ff54a471..1e905e732 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 1bab492ea..40f34dba7 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 1547b9424..b0cfd20a1 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 aec1cf748..52c86a856 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 b95117460..755e82e3c 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 ae8d1d8b3..fd49b7a82 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 5ff9d408c..4dd934945 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 f2e0583d0..eb41a6940 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 30a0f48dd..3a27a6b70 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 8e39a7326..5b5a6311d 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 4d9a39bb4..741ac5306 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 e073fd578..0951ced54 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 6f0f8d8f0..15b2cbb24 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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 a8b663f31..425d04ec8 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 @@ -33,8 +33,8 @@ namespace cuvs::neighbors::cagra::detail::single_cta_search { TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T, SAMPLE_FILTER_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/refine_host_float_float.cpp b/cpp/src/neighbors/detail/refine_host_float_float.cpp index cd4790f94..ceea544a0 100644 --- a/cpp/src/neighbors/detail/refine_host_float_float.cpp +++ b/cpp/src/neighbors/detail/refine_host_float_float.cpp @@ -17,11 +17,11 @@ #define instantiate_raft_neighbors_refine(IdxT, DataT, DistanceT, ExtentsT) \ template void cuvs::neighbors::detail::refine_host( \ - raft::host_matrix_view dataset, \ - raft::host_matrix_view queries, \ - raft::host_matrix_view neighbor_candidates, \ - raft::host_matrix_view indices, \ - raft::host_matrix_view distances, \ + raft::host_matrix_view dataset, \ + raft::host_matrix_view queries, \ + raft::host_matrix_view neighbor_candidates, \ + raft::host_matrix_view indices, \ + raft::host_matrix_view distances, \ distance::DistanceType metric); instantiate_raft_neighbors_refine(int64_t, float, float, int64_t); diff --git a/cpp/src/neighbors/detail/refine_host_int8_t_float.cpp b/cpp/src/neighbors/detail/refine_host_int8_t_float.cpp index 2511ba7cc..c236740b4 100644 --- a/cpp/src/neighbors/detail/refine_host_int8_t_float.cpp +++ b/cpp/src/neighbors/detail/refine_host_int8_t_float.cpp @@ -18,11 +18,11 @@ #define instantiate_raft_neighbors_refine(IdxT, DataT, DistanceT, ExtentsT) \ template void cuvs::neighbors::detail::refine_host( \ - raft::host_matrix_view dataset, \ - raft::host_matrix_view queries, \ - raft::host_matrix_view neighbor_candidates, \ - raft::host_matrix_view indices, \ - raft::host_matrix_view distances, \ + raft::host_matrix_view dataset, \ + raft::host_matrix_view queries, \ + raft::host_matrix_view neighbor_candidates, \ + raft::host_matrix_view indices, \ + raft::host_matrix_view distances, \ distance::DistanceType metric); instantiate_raft_neighbors_refine(int64_t, int8_t, float, int64_t); diff --git a/cpp/src/neighbors/detail/refine_host_uint8_t_float.cpp b/cpp/src/neighbors/detail/refine_host_uint8_t_float.cpp index 98367faad..5aee0f29e 100644 --- a/cpp/src/neighbors/detail/refine_host_uint8_t_float.cpp +++ b/cpp/src/neighbors/detail/refine_host_uint8_t_float.cpp @@ -18,11 +18,11 @@ #define instantiate_raft_neighbors_refine(IdxT, DataT, DistanceT, ExtentsT) \ template void cuvs::neighbors::detail::refine_host( \ - raft::host_matrix_view dataset, \ - raft::host_matrix_view queries, \ - raft::host_matrix_view neighbor_candidates, \ - raft::host_matrix_view indices, \ - raft::host_matrix_view distances, \ + raft::host_matrix_view dataset, \ + raft::host_matrix_view queries, \ + raft::host_matrix_view neighbor_candidates, \ + raft::host_matrix_view indices, \ + raft::host_matrix_view distances, \ distance::DistanceType metric); instantiate_raft_neighbors_refine(int64_t, uint8_t, float, int64_t); diff --git a/cpp/src/neighbors/ivf_flat_00_generate.py b/cpp/src/neighbors/ivf_flat_00_generate.py index a83981572..dc7d7374f 100644 --- a/cpp/src/neighbors/ivf_flat_00_generate.py +++ b/cpp/src/neighbors/ivf_flat_00_generate.py @@ -59,13 +59,13 @@ template auto cuvs::neighbors::ivf_flat::build( \\ raft::resources const& handle, \\ const cuvs::neighbors::ivf_flat::index_params& params, \\ - raft::device_matrix_view dataset) \\ + raft::device_matrix_view dataset) \\ ->cuvs::neighbors::ivf_flat::index; \\ \\ template void cuvs::neighbors::ivf_flat::build( \\ raft::resources const& handle, \\ const cuvs::neighbors::ivf_flat::index_params& params, \\ - raft::device_matrix_view dataset, \\ + raft::device_matrix_view dataset, \\ cuvs::neighbors::ivf_flat::index& idx); """ @@ -81,7 +81,7 @@ \\ template auto cuvs::neighbors::ivf_flat::extend( \\ raft::resources const& handle, \\ - raft::device_matrix_view new_vectors, \\ + raft::device_matrix_view new_vectors, \\ std::optional> new_indices, \\ const cuvs::neighbors::ivf_flat::index& orig_index) \\ ->cuvs::neighbors::ivf_flat::index; \\ @@ -95,7 +95,7 @@ \\ template void cuvs::neighbors::ivf_flat::extend( \\ raft::resources const& handle, \\ - raft::device_matrix_view new_vectors, \\ + raft::device_matrix_view new_vectors, \\ std::optional> new_indices, \\ cuvs::neighbors::ivf_flat::index* index); """ @@ -117,9 +117,9 @@ 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); + raft::device_matrix_view queries, \\ + raft::device_matrix_view neighbors, \\ + raft::device_matrix_view distances); """ macros = dict( diff --git a/cpp/src/neighbors/ivf_flat_build_float_int64_t.cu b/cpp/src/neighbors/ivf_flat_build_float_int64_t.cu index 8a3c5dfa8..794e435bd 100644 --- a/cpp/src/neighbors/ivf_flat_build_float_int64_t.cu +++ b/cpp/src/neighbors/ivf_flat_build_float_int64_t.cu @@ -25,25 +25,25 @@ #include -#define instantiate_raft_neighbors_ivf_flat_build(T, IdxT) \ - 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; \ - \ - 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; \ - \ - template void cuvs::neighbors::ivf_flat::build( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_flat::index_params& params, \ - raft::device_matrix_view dataset, \ +#define instantiate_raft_neighbors_ivf_flat_build(T, IdxT) \ + 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; \ + \ + 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; \ + \ + 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); diff --git a/cpp/src/neighbors/ivf_flat_build_int8_t_int64_t.cu b/cpp/src/neighbors/ivf_flat_build_int8_t_int64_t.cu index 90232f5b6..a2e9bd828 100644 --- a/cpp/src/neighbors/ivf_flat_build_int8_t_int64_t.cu +++ b/cpp/src/neighbors/ivf_flat_build_int8_t_int64_t.cu @@ -25,25 +25,25 @@ #include -#define instantiate_raft_neighbors_ivf_flat_build(T, IdxT) \ - 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; \ - \ - 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; \ - \ - template void cuvs::neighbors::ivf_flat::build( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_flat::index_params& params, \ - raft::device_matrix_view dataset, \ +#define instantiate_raft_neighbors_ivf_flat_build(T, IdxT) \ + 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; \ + \ + 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; \ + \ + 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(int8_t, int64_t); diff --git a/cpp/src/neighbors/ivf_flat_build_uint8_t_int64_t.cu b/cpp/src/neighbors/ivf_flat_build_uint8_t_int64_t.cu index 8ad75bb94..51d855374 100644 --- a/cpp/src/neighbors/ivf_flat_build_uint8_t_int64_t.cu +++ b/cpp/src/neighbors/ivf_flat_build_uint8_t_int64_t.cu @@ -25,25 +25,25 @@ #include -#define instantiate_raft_neighbors_ivf_flat_build(T, IdxT) \ - 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; \ - \ - 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; \ - \ - template void cuvs::neighbors::ivf_flat::build( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_flat::index_params& params, \ - raft::device_matrix_view dataset, \ +#define instantiate_raft_neighbors_ivf_flat_build(T, IdxT) \ + 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; \ + \ + 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; \ + \ + 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(uint8_t, int64_t); diff --git a/cpp/src/neighbors/ivf_flat_extend_float_int64_t.cu b/cpp/src/neighbors/ivf_flat_extend_float_int64_t.cu index 4b898390d..2e825938c 100644 --- a/cpp/src/neighbors/ivf_flat_extend_float_int64_t.cu +++ b/cpp/src/neighbors/ivf_flat_extend_float_int64_t.cu @@ -36,7 +36,7 @@ \ template auto cuvs::neighbors::ivf_flat::extend( \ raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ + raft::device_matrix_view new_vectors, \ std::optional> new_indices, \ const cuvs::neighbors::ivf_flat::index& orig_index) \ ->cuvs::neighbors::ivf_flat::index; \ @@ -50,7 +50,7 @@ \ template void cuvs::neighbors::ivf_flat::extend( \ raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ + 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); diff --git a/cpp/src/neighbors/ivf_flat_extend_int8_t_int64_t.cu b/cpp/src/neighbors/ivf_flat_extend_int8_t_int64_t.cu index c618a1297..fa11b4472 100644 --- a/cpp/src/neighbors/ivf_flat_extend_int8_t_int64_t.cu +++ b/cpp/src/neighbors/ivf_flat_extend_int8_t_int64_t.cu @@ -36,7 +36,7 @@ \ template auto cuvs::neighbors::ivf_flat::extend( \ raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ + raft::device_matrix_view new_vectors, \ std::optional> new_indices, \ const cuvs::neighbors::ivf_flat::index& orig_index) \ ->cuvs::neighbors::ivf_flat::index; \ @@ -50,7 +50,7 @@ \ template void cuvs::neighbors::ivf_flat::extend( \ raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ + raft::device_matrix_view new_vectors, \ std::optional> new_indices, \ cuvs::neighbors::ivf_flat::index* index); instantiate_raft_neighbors_ivf_flat_extend(int8_t, int64_t); diff --git a/cpp/src/neighbors/ivf_flat_extend_uint8_t_int64_t.cu b/cpp/src/neighbors/ivf_flat_extend_uint8_t_int64_t.cu index 25fd67cf5..8e4e24349 100644 --- a/cpp/src/neighbors/ivf_flat_extend_uint8_t_int64_t.cu +++ b/cpp/src/neighbors/ivf_flat_extend_uint8_t_int64_t.cu @@ -36,7 +36,7 @@ \ template auto cuvs::neighbors::ivf_flat::extend( \ raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ + raft::device_matrix_view new_vectors, \ std::optional> new_indices, \ const cuvs::neighbors::ivf_flat::index& orig_index) \ ->cuvs::neighbors::ivf_flat::index; \ @@ -50,7 +50,7 @@ \ template void cuvs::neighbors::ivf_flat::extend( \ raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ + raft::device_matrix_view new_vectors, \ std::optional> new_indices, \ cuvs::neighbors::ivf_flat::index* index); instantiate_raft_neighbors_ivf_flat_extend(uint8_t, int64_t); diff --git a/cpp/src/neighbors/ivf_flat_search_float_int64_t.cu b/cpp/src/neighbors/ivf_flat_search_float_int64_t.cu index 21318072b..d0072c0ff 100644 --- a/cpp/src/neighbors/ivf_flat_search_float_int64_t.cu +++ b/cpp/src/neighbors/ivf_flat_search_float_int64_t.cu @@ -25,25 +25,25 @@ #include -#define instantiate_raft_neighbors_ivf_flat_search(T, IdxT) \ - 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); \ - \ - 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); +#define instantiate_raft_neighbors_ivf_flat_search(T, IdxT) \ + 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); \ + \ + 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); #undef instantiate_raft_neighbors_ivf_flat_search diff --git a/cpp/src/neighbors/ivf_flat_search_int8_t_int64_t.cu b/cpp/src/neighbors/ivf_flat_search_int8_t_int64_t.cu index 6f0eea00c..ecdc2a2ed 100644 --- a/cpp/src/neighbors/ivf_flat_search_int8_t_int64_t.cu +++ b/cpp/src/neighbors/ivf_flat_search_int8_t_int64_t.cu @@ -25,25 +25,25 @@ #include -#define instantiate_raft_neighbors_ivf_flat_search(T, IdxT) \ - 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); \ - \ - 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); +#define instantiate_raft_neighbors_ivf_flat_search(T, IdxT) \ + 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); \ + \ + 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(int8_t, int64_t); #undef instantiate_raft_neighbors_ivf_flat_search diff --git a/cpp/src/neighbors/ivf_flat_search_uint8_t_int64_t.cu b/cpp/src/neighbors/ivf_flat_search_uint8_t_int64_t.cu index 7f7550359..a87486004 100644 --- a/cpp/src/neighbors/ivf_flat_search_uint8_t_int64_t.cu +++ b/cpp/src/neighbors/ivf_flat_search_uint8_t_int64_t.cu @@ -25,25 +25,25 @@ #include -#define instantiate_raft_neighbors_ivf_flat_search(T, IdxT) \ - 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); \ - \ - 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); +#define instantiate_raft_neighbors_ivf_flat_search(T, IdxT) \ + 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); \ + \ + 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(uint8_t, int64_t); #undef instantiate_raft_neighbors_ivf_flat_search diff --git a/cpp/src/neighbors/ivfpq_build_float_int64_t.cu b/cpp/src/neighbors/ivfpq_build_float_int64_t.cu index ec844fe50..672bbe732 100644 --- a/cpp/src/neighbors/ivfpq_build_float_int64_t.cu +++ b/cpp/src/neighbors/ivfpq_build_float_int64_t.cu @@ -21,7 +21,7 @@ template cuvs::neighbors::ivf_pq::index cuvs::neighbors::ivf_pq::build( \ raft::resources const& handle, \ const cuvs::neighbors::ivf_pq::index_params& params, \ - raft::device_matrix_view dataset); \ + raft::device_matrix_view dataset); \ \ template auto cuvs::neighbors::ivf_pq::build( \ raft::resources const& handle, \ diff --git a/cpp/src/neighbors/ivfpq_build_int8_t_int64_t.cu b/cpp/src/neighbors/ivfpq_build_int8_t_int64_t.cu index ab34c6ef7..f3d3a4e5f 100644 --- a/cpp/src/neighbors/ivfpq_build_int8_t_int64_t.cu +++ b/cpp/src/neighbors/ivfpq_build_int8_t_int64_t.cu @@ -21,7 +21,7 @@ template cuvs::neighbors::ivf_pq::index cuvs::neighbors::ivf_pq::build( \ raft::resources const& handle, \ const cuvs::neighbors::ivf_pq::index_params& params, \ - raft::device_matrix_view dataset); \ + raft::device_matrix_view dataset); \ \ template auto cuvs::neighbors::ivf_pq::build( \ raft::resources const& handle, \ diff --git a/cpp/src/neighbors/ivfpq_build_uint8_t_int64_t.cu b/cpp/src/neighbors/ivfpq_build_uint8_t_int64_t.cu index e6348fbff..ffd630b1e 100644 --- a/cpp/src/neighbors/ivfpq_build_uint8_t_int64_t.cu +++ b/cpp/src/neighbors/ivfpq_build_uint8_t_int64_t.cu @@ -21,7 +21,7 @@ template cuvs::neighbors::ivf_pq::index cuvs::neighbors::ivf_pq::build( \ raft::resources const& handle, \ const cuvs::neighbors::ivf_pq::index_params& params, \ - raft::device_matrix_view dataset); \ + raft::device_matrix_view dataset); \ \ template auto cuvs::neighbors::ivf_pq::build( \ raft::resources const& handle, \ diff --git a/cpp/src/neighbors/ivfpq_extend_float_int64_t.cu b/cpp/src/neighbors/ivfpq_extend_float_int64_t.cu index 42c024891..59385c8b7 100644 --- a/cpp/src/neighbors/ivfpq_extend_float_int64_t.cu +++ b/cpp/src/neighbors/ivfpq_extend_float_int64_t.cu @@ -17,32 +17,32 @@ #include #include // cuvs::neighbors::ivf_pq::index -#define instantiate_raft_neighbors_ivf_pq_extend(T, IdxT) \ - template cuvs::neighbors::ivf_pq::index cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ - std::optional> new_indices, \ - const cuvs::neighbors::ivf_pq::index& idx); \ - \ - template void cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ - std::optional> new_indices, \ - cuvs::neighbors::ivf_pq::index* idx); \ - \ - template auto cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - const T* new_vectors, \ - const IdxT* new_indices, \ - IdxT n_rows) \ - ->cuvs::neighbors::ivf_pq::index; \ - \ - template void cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - cuvs::neighbors::ivf_pq::index* idx, \ - const T* new_vectors, \ - const IdxT* new_indices, \ +#define instantiate_raft_neighbors_ivf_pq_extend(T, IdxT) \ + template cuvs::neighbors::ivf_pq::index cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + raft::device_matrix_view new_vectors, \ + std::optional> new_indices, \ + const cuvs::neighbors::ivf_pq::index& idx); \ + \ + template void cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + raft::device_matrix_view new_vectors, \ + std::optional> new_indices, \ + cuvs::neighbors::ivf_pq::index* idx); \ + \ + template auto cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + const T* new_vectors, \ + const IdxT* new_indices, \ + IdxT n_rows) \ + ->cuvs::neighbors::ivf_pq::index; \ + \ + template void cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + cuvs::neighbors::ivf_pq::index* idx, \ + const T* new_vectors, \ + const IdxT* new_indices, \ IdxT n_rows); instantiate_raft_neighbors_ivf_pq_extend(float, int64_t); diff --git a/cpp/src/neighbors/ivfpq_extend_int8_t_int64_t.cu b/cpp/src/neighbors/ivfpq_extend_int8_t_int64_t.cu index 5fc0fc9ed..7fad247d1 100644 --- a/cpp/src/neighbors/ivfpq_extend_int8_t_int64_t.cu +++ b/cpp/src/neighbors/ivfpq_extend_int8_t_int64_t.cu @@ -17,32 +17,32 @@ #include #include // cuvs::neighbors::ivf_pq::index -#define instantiate_raft_neighbors_ivf_pq_extend(T, IdxT) \ - template cuvs::neighbors::ivf_pq::index cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ - std::optional> new_indices, \ - const cuvs::neighbors::ivf_pq::index& idx); \ - \ - template void cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ - std::optional> new_indices, \ - cuvs::neighbors::ivf_pq::index* idx); \ - \ - template auto cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - const T* new_vectors, \ - const IdxT* new_indices, \ - IdxT n_rows) \ - ->cuvs::neighbors::ivf_pq::index; \ - \ - template void cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - cuvs::neighbors::ivf_pq::index* idx, \ - const T* new_vectors, \ - const IdxT* new_indices, \ +#define instantiate_raft_neighbors_ivf_pq_extend(T, IdxT) \ + template cuvs::neighbors::ivf_pq::index cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + raft::device_matrix_view new_vectors, \ + std::optional> new_indices, \ + const cuvs::neighbors::ivf_pq::index& idx); \ + \ + template void cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + raft::device_matrix_view new_vectors, \ + std::optional> new_indices, \ + cuvs::neighbors::ivf_pq::index* idx); \ + \ + template auto cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + const T* new_vectors, \ + const IdxT* new_indices, \ + IdxT n_rows) \ + ->cuvs::neighbors::ivf_pq::index; \ + \ + template void cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + cuvs::neighbors::ivf_pq::index* idx, \ + const T* new_vectors, \ + const IdxT* new_indices, \ IdxT n_rows); instantiate_raft_neighbors_ivf_pq_extend(int8_t, int64_t); diff --git a/cpp/src/neighbors/ivfpq_extend_uint8_t_int64_t.cu b/cpp/src/neighbors/ivfpq_extend_uint8_t_int64_t.cu index 8b3c3fa03..7bbdd5ffd 100644 --- a/cpp/src/neighbors/ivfpq_extend_uint8_t_int64_t.cu +++ b/cpp/src/neighbors/ivfpq_extend_uint8_t_int64_t.cu @@ -17,32 +17,32 @@ #include #include // cuvs::neighbors::ivf_pq::index -#define instantiate_raft_neighbors_ivf_pq_extend(T, IdxT) \ - template cuvs::neighbors::ivf_pq::index cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ - std::optional> new_indices, \ - const cuvs::neighbors::ivf_pq::index& idx); \ - \ - template void cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - raft::device_matrix_view new_vectors, \ - std::optional> new_indices, \ - cuvs::neighbors::ivf_pq::index* idx); \ - \ - template auto cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - const T* new_vectors, \ - const IdxT* new_indices, \ - IdxT n_rows) \ - ->cuvs::neighbors::ivf_pq::index; \ - \ - template void cuvs::neighbors::ivf_pq::extend( \ - raft::resources const& handle, \ - cuvs::neighbors::ivf_pq::index* idx, \ - const T* new_vectors, \ - const IdxT* new_indices, \ +#define instantiate_raft_neighbors_ivf_pq_extend(T, IdxT) \ + template cuvs::neighbors::ivf_pq::index cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + raft::device_matrix_view new_vectors, \ + std::optional> new_indices, \ + const cuvs::neighbors::ivf_pq::index& idx); \ + \ + template void cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + raft::device_matrix_view new_vectors, \ + std::optional> new_indices, \ + cuvs::neighbors::ivf_pq::index* idx); \ + \ + template auto cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + const T* new_vectors, \ + const IdxT* new_indices, \ + IdxT n_rows) \ + ->cuvs::neighbors::ivf_pq::index; \ + \ + template void cuvs::neighbors::ivf_pq::extend( \ + raft::resources const& handle, \ + cuvs::neighbors::ivf_pq::index* idx, \ + const T* new_vectors, \ + const IdxT* new_indices, \ IdxT n_rows); instantiate_raft_neighbors_ivf_pq_extend(uint8_t, int64_t); diff --git a/cpp/src/neighbors/ivfpq_search_float_int64_t.cu b/cpp/src/neighbors/ivfpq_search_float_int64_t.cu index c32a40ad2..31ce6e8df 100644 --- a/cpp/src/neighbors/ivfpq_search_float_int64_t.cu +++ b/cpp/src/neighbors/ivfpq_search_float_int64_t.cu @@ -17,24 +17,24 @@ #include #include // cuvs::neighbors::ivf_pq::index -#define instantiate_raft_neighbors_ivf_pq_search(T, IdxT) \ - template void cuvs::neighbors::ivf_pq::search( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::search_params& params, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - raft::device_matrix_view queries, \ - raft::device_matrix_view neighbors, \ - raft::device_matrix_view distances); \ - \ - template void cuvs::neighbors::ivf_pq::search( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::search_params& params, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - const T* queries, \ - uint32_t n_queries, \ - uint32_t k, \ - IdxT* neighbors, \ - float* distances, \ +#define instantiate_raft_neighbors_ivf_pq_search(T, IdxT) \ + template void cuvs::neighbors::ivf_pq::search( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::search_params& params, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + raft::device_matrix_view queries, \ + raft::device_matrix_view neighbors, \ + raft::device_matrix_view distances); \ + \ + template void cuvs::neighbors::ivf_pq::search( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::search_params& params, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + const T* queries, \ + uint32_t n_queries, \ + uint32_t k, \ + IdxT* neighbors, \ + float* distances, \ rmm::mr::device_memory_resource* mr) instantiate_raft_neighbors_ivf_pq_search(float, int64_t); diff --git a/cpp/src/neighbors/ivfpq_search_int8_t_int64_t.cu b/cpp/src/neighbors/ivfpq_search_int8_t_int64_t.cu index 6a2e59fbf..5b50b3b19 100644 --- a/cpp/src/neighbors/ivfpq_search_int8_t_int64_t.cu +++ b/cpp/src/neighbors/ivfpq_search_int8_t_int64_t.cu @@ -17,24 +17,24 @@ #include #include // cuvs::neighbors::ivf_pq::index -#define instantiate_raft_neighbors_ivf_pq_search(T, IdxT) \ - template void cuvs::neighbors::ivf_pq::search( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::search_params& params, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - raft::device_matrix_view queries, \ - raft::device_matrix_view neighbors, \ - raft::device_matrix_view distances); \ - \ - template void cuvs::neighbors::ivf_pq::search( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::search_params& params, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - const T* queries, \ - uint32_t n_queries, \ - uint32_t k, \ - IdxT* neighbors, \ - float* distances, \ +#define instantiate_raft_neighbors_ivf_pq_search(T, IdxT) \ + template void cuvs::neighbors::ivf_pq::search( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::search_params& params, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + raft::device_matrix_view queries, \ + raft::device_matrix_view neighbors, \ + raft::device_matrix_view distances); \ + \ + template void cuvs::neighbors::ivf_pq::search( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::search_params& params, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + const T* queries, \ + uint32_t n_queries, \ + uint32_t k, \ + IdxT* neighbors, \ + float* distances, \ rmm::mr::device_memory_resource* mr) instantiate_raft_neighbors_ivf_pq_search(int8_t, int64_t); diff --git a/cpp/src/neighbors/ivfpq_search_uint8_t_int64_t.cu b/cpp/src/neighbors/ivfpq_search_uint8_t_int64_t.cu index ae7246c72..100c3b49e 100644 --- a/cpp/src/neighbors/ivfpq_search_uint8_t_int64_t.cu +++ b/cpp/src/neighbors/ivfpq_search_uint8_t_int64_t.cu @@ -17,24 +17,24 @@ #include #include // cuvs::neighbors::ivf_pq::index -#define instantiate_raft_neighbors_ivf_pq_search(T, IdxT) \ - template void cuvs::neighbors::ivf_pq::search( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::search_params& params, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - raft::device_matrix_view queries, \ - raft::device_matrix_view neighbors, \ - raft::device_matrix_view distances); \ - \ - template void cuvs::neighbors::ivf_pq::search( \ - raft::resources const& handle, \ - const cuvs::neighbors::ivf_pq::search_params& params, \ - const cuvs::neighbors::ivf_pq::index& idx, \ - const T* queries, \ - uint32_t n_queries, \ - uint32_t k, \ - IdxT* neighbors, \ - float* distances, \ +#define instantiate_raft_neighbors_ivf_pq_search(T, IdxT) \ + template void cuvs::neighbors::ivf_pq::search( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::search_params& params, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + raft::device_matrix_view queries, \ + raft::device_matrix_view neighbors, \ + raft::device_matrix_view distances); \ + \ + template void cuvs::neighbors::ivf_pq::search( \ + raft::resources const& handle, \ + const cuvs::neighbors::ivf_pq::search_params& params, \ + const cuvs::neighbors::ivf_pq::index& idx, \ + const T* queries, \ + uint32_t n_queries, \ + uint32_t k, \ + IdxT* neighbors, \ + float* distances, \ rmm::mr::device_memory_resource* mr) instantiate_raft_neighbors_ivf_pq_search(uint8_t, int64_t); diff --git a/cpp/src/neighbors/refine_00_generate.py b/cpp/src/neighbors/refine_00_generate.py index de07c1215..45cfc473e 100644 --- a/cpp/src/neighbors/refine_00_generate.py +++ b/cpp/src/neighbors/refine_00_generate.py @@ -43,20 +43,20 @@ #define instantiate_raft_neighbors_refine(idx_t, data_t, distance_t, matrix_idx) \\ template void cuvs::neighbors::refine( \\ raft::resources const& handle, \\ - raft::device_matrix_view dataset, \\ - raft::device_matrix_view queries, \\ - raft::device_matrix_view neighbor_candidates, \\ - raft::device_matrix_view indices, \\ - raft::device_matrix_view distances, \\ + raft::device_matrix_view dataset, \\ + raft::device_matrix_view queries, \\ + raft::device_matrix_view neighbor_candidates, \\ + raft::device_matrix_view indices, \\ + raft::device_matrix_view distances, \\ cuvs::distance::DistanceType metric); \\ \\ template void cuvs::neighbors::refine( \\ raft::resources const& handle, \\ - raft::host_matrix_view dataset, \\ - raft::host_matrix_view queries, \\ - raft::host_matrix_view neighbor_candidates, \\ - raft::host_matrix_view indices, \\ - raft::host_matrix_view distances, \\ + raft::host_matrix_view dataset, \\ + raft::host_matrix_view queries, \\ + raft::host_matrix_view neighbor_candidates, \\ + raft::host_matrix_view indices, \\ + raft::host_matrix_view distances, \\ cuvs::distance::DistanceType metric); """ diff --git a/cpp/src/neighbors/refine_float_float.cu b/cpp/src/neighbors/refine_float_float.cu index 4c33311d6..39d51c1bc 100644 --- a/cpp/src/neighbors/refine_float_float.cu +++ b/cpp/src/neighbors/refine_float_float.cu @@ -26,23 +26,23 @@ #include -#define instantiate_raft_neighbors_refine(idx_t, data_t, distance_t, matrix_idx) \ - template void cuvs::neighbors::refine( \ - raft::resources const& handle, \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view queries, \ - raft::device_matrix_view neighbor_candidates, \ - raft::device_matrix_view indices, \ - raft::device_matrix_view distances, \ - cuvs::distance::DistanceType metric); \ - \ - template void cuvs::neighbors::refine( \ - raft::resources const& handle, \ - raft::host_matrix_view dataset, \ - raft::host_matrix_view queries, \ - raft::host_matrix_view neighbor_candidates, \ - raft::host_matrix_view indices, \ - raft::host_matrix_view distances, \ +#define instantiate_raft_neighbors_refine(idx_t, data_t, distance_t, matrix_idx) \ + template void cuvs::neighbors::refine( \ + raft::resources const& handle, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view queries, \ + raft::device_matrix_view neighbor_candidates, \ + raft::device_matrix_view indices, \ + raft::device_matrix_view distances, \ + cuvs::distance::DistanceType metric); \ + \ + template void cuvs::neighbors::refine( \ + raft::resources const& handle, \ + raft::host_matrix_view dataset, \ + raft::host_matrix_view queries, \ + raft::host_matrix_view neighbor_candidates, \ + raft::host_matrix_view indices, \ + raft::host_matrix_view distances, \ cuvs::distance::DistanceType metric); instantiate_raft_neighbors_refine(int64_t, float, float, int64_t); diff --git a/cpp/src/neighbors/refine_int8_t_float.cu b/cpp/src/neighbors/refine_int8_t_float.cu index 6f394e577..f65590a4f 100644 --- a/cpp/src/neighbors/refine_int8_t_float.cu +++ b/cpp/src/neighbors/refine_int8_t_float.cu @@ -26,23 +26,23 @@ #include -#define instantiate_raft_neighbors_refine(idx_t, data_t, distance_t, matrix_idx) \ - template void cuvs::neighbors::refine( \ - raft::resources const& handle, \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view queries, \ - raft::device_matrix_view neighbor_candidates, \ - raft::device_matrix_view indices, \ - raft::device_matrix_view distances, \ - cuvs::distance::DistanceType metric); \ - \ - template void cuvs::neighbors::refine( \ - raft::resources const& handle, \ - raft::host_matrix_view dataset, \ - raft::host_matrix_view queries, \ - raft::host_matrix_view neighbor_candidates, \ - raft::host_matrix_view indices, \ - raft::host_matrix_view distances, \ +#define instantiate_raft_neighbors_refine(idx_t, data_t, distance_t, matrix_idx) \ + template void cuvs::neighbors::refine( \ + raft::resources const& handle, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view queries, \ + raft::device_matrix_view neighbor_candidates, \ + raft::device_matrix_view indices, \ + raft::device_matrix_view distances, \ + cuvs::distance::DistanceType metric); \ + \ + template void cuvs::neighbors::refine( \ + raft::resources const& handle, \ + raft::host_matrix_view dataset, \ + raft::host_matrix_view queries, \ + raft::host_matrix_view neighbor_candidates, \ + raft::host_matrix_view indices, \ + raft::host_matrix_view distances, \ cuvs::distance::DistanceType metric); instantiate_raft_neighbors_refine(int64_t, int8_t, float, int64_t); diff --git a/cpp/src/neighbors/refine_uint8_t_float.cu b/cpp/src/neighbors/refine_uint8_t_float.cu index 527458e71..ba76e4904 100644 --- a/cpp/src/neighbors/refine_uint8_t_float.cu +++ b/cpp/src/neighbors/refine_uint8_t_float.cu @@ -26,23 +26,23 @@ #include -#define instantiate_raft_neighbors_refine(idx_t, data_t, distance_t, matrix_idx) \ - template void cuvs::neighbors::refine( \ - raft::resources const& handle, \ - raft::device_matrix_view dataset, \ - raft::device_matrix_view queries, \ - raft::device_matrix_view neighbor_candidates, \ - raft::device_matrix_view indices, \ - raft::device_matrix_view distances, \ - cuvs::distance::DistanceType metric); \ - \ - template void cuvs::neighbors::refine( \ - raft::resources const& handle, \ - raft::host_matrix_view dataset, \ - raft::host_matrix_view queries, \ - raft::host_matrix_view neighbor_candidates, \ - raft::host_matrix_view indices, \ - raft::host_matrix_view distances, \ +#define instantiate_raft_neighbors_refine(idx_t, data_t, distance_t, matrix_idx) \ + template void cuvs::neighbors::refine( \ + raft::resources const& handle, \ + raft::device_matrix_view dataset, \ + raft::device_matrix_view queries, \ + raft::device_matrix_view neighbor_candidates, \ + raft::device_matrix_view indices, \ + raft::device_matrix_view distances, \ + cuvs::distance::DistanceType metric); \ + \ + template void cuvs::neighbors::refine( \ + raft::resources const& handle, \ + raft::host_matrix_view dataset, \ + raft::host_matrix_view queries, \ + raft::host_matrix_view neighbor_candidates, \ + raft::host_matrix_view indices, \ + raft::host_matrix_view distances, \ cuvs::distance::DistanceType metric); instantiate_raft_neighbors_refine(int64_t, uint8_t, float, int64_t);