From a9fc649ca79f60b60d9d171edbbb39fec300a5af Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Mon, 24 Jul 2023 13:18:16 -0400 Subject: [PATCH 1/3] Promoting CAGRA from experimental. Follow-on commit will contain some documentation around it --- cpp/bench/ann/src/raft/raft_cagra_wrapper.h | 16 ++++++++-------- cpp/bench/prims/neighbors/cagra_bench.cuh | 8 ++++---- cpp/include/raft/neighbors/cagra.cuh | 18 +++++++++++++----- cpp/include/raft/neighbors/cagra_serialize.cuh | 14 ++++++++++---- cpp/include/raft/neighbors/cagra_types.hpp | 10 +++++++++- .../raft/neighbors/detail/cagra/bitonic.hpp | 4 ++-- .../neighbors/detail/cagra/cagra_build.cuh | 4 ++-- .../neighbors/detail/cagra/cagra_search.cuh | 4 ++-- .../neighbors/detail/cagra/cagra_serialize.cuh | 4 ++-- .../detail/cagra/compute_distance.hpp | 4 ++-- .../neighbors/detail/cagra/device_common.hpp | 4 ++-- .../raft/neighbors/detail/cagra/factory.cuh | 4 ++-- .../raft/neighbors/detail/cagra/fragment.hpp | 4 ++-- .../raft/neighbors/detail/cagra/graph_core.cuh | 4 ++-- .../raft/neighbors/detail/cagra/hashmap.hpp | 4 ++-- .../detail/cagra/search_multi_cta.cuh | 4 ++-- .../cagra/search_multi_cta_kernel-ext.cuh | 4 ++-- .../cagra/search_multi_cta_kernel-inl.cuh | 4 ++-- .../detail/cagra/search_multi_kernel.cuh | 4 ++-- .../neighbors/detail/cagra/search_plan.cuh | 4 ++-- .../detail/cagra/search_single_cta.cuh | 4 ++-- .../cagra/search_single_cta_kernel-ext.cuh | 4 ++-- .../cagra/search_single_cta_kernel-inl.cuh | 4 ++-- .../neighbors/detail/cagra/topk_by_radix.cuh | 4 ++-- .../detail/cagra/topk_for_cagra/topk.h | 4 ++-- .../detail/cagra/topk_for_cagra/topk_core.cuh | 4 ++-- .../raft/neighbors/detail/cagra/utils.hpp | 4 ++-- .../cagra/search_multi_cta_00_generate.py | 4 ++-- ...earch_multi_cta_float_uint32_dim1024_t32.cu | 4 ++-- .../search_multi_cta_float_uint32_dim128_t8.cu | 4 ++-- ...search_multi_cta_float_uint32_dim256_t16.cu | 4 ++-- ...search_multi_cta_float_uint32_dim512_t32.cu | 4 ++-- ...earch_multi_cta_float_uint64_dim1024_t32.cu | 4 ++-- .../search_multi_cta_float_uint64_dim128_t8.cu | 4 ++-- ...search_multi_cta_float_uint64_dim256_t16.cu | 4 ++-- ...search_multi_cta_float_uint64_dim512_t32.cu | 4 ++-- ...search_multi_cta_int8_uint32_dim1024_t32.cu | 4 ++-- .../search_multi_cta_int8_uint32_dim128_t8.cu | 4 ++-- .../search_multi_cta_int8_uint32_dim256_t16.cu | 4 ++-- .../search_multi_cta_int8_uint32_dim512_t32.cu | 4 ++-- ...earch_multi_cta_uint8_uint32_dim1024_t32.cu | 4 ++-- .../search_multi_cta_uint8_uint32_dim128_t8.cu | 4 ++-- ...search_multi_cta_uint8_uint32_dim256_t16.cu | 4 ++-- ...search_multi_cta_uint8_uint32_dim512_t32.cu | 4 ++-- .../cagra/search_single_cta_00_generate.py | 4 ++-- ...arch_single_cta_float_uint32_dim1024_t32.cu | 4 ++-- ...search_single_cta_float_uint32_dim128_t8.cu | 4 ++-- ...earch_single_cta_float_uint32_dim256_t16.cu | 4 ++-- ...earch_single_cta_float_uint32_dim512_t32.cu | 4 ++-- ...arch_single_cta_float_uint64_dim1024_t32.cu | 4 ++-- ...search_single_cta_float_uint64_dim128_t8.cu | 4 ++-- ...earch_single_cta_float_uint64_dim256_t16.cu | 4 ++-- ...earch_single_cta_float_uint64_dim512_t32.cu | 4 ++-- ...earch_single_cta_int8_uint32_dim1024_t32.cu | 4 ++-- .../search_single_cta_int8_uint32_dim128_t8.cu | 4 ++-- ...search_single_cta_int8_uint32_dim256_t16.cu | 4 ++-- ...search_single_cta_int8_uint32_dim512_t32.cu | 4 ++-- ...arch_single_cta_uint8_uint32_dim1024_t32.cu | 4 ++-- ...search_single_cta_uint8_uint32_dim128_t8.cu | 4 ++-- ...earch_single_cta_uint8_uint32_dim256_t16.cu | 4 ++-- ...earch_single_cta_uint8_uint32_dim512_t32.cu | 4 ++-- cpp/test/neighbors/ann_cagra.cuh | 4 ++-- .../ann_cagra/search_kernel_uint64_t.cuh | 4 ++-- docs/source/cpp_api/neighbors_cagra.rst | 2 +- 64 files changed, 161 insertions(+), 139 deletions(-) diff --git a/cpp/bench/ann/src/raft/raft_cagra_wrapper.h b/cpp/bench/ann/src/raft/raft_cagra_wrapper.h index 399fd6a0a8..aa3464e2b9 100644 --- a/cpp/bench/ann/src/raft/raft_cagra_wrapper.h +++ b/cpp/bench/ann/src/raft/raft_cagra_wrapper.h @@ -50,7 +50,7 @@ class RaftCagra : public ANN { unsigned itopk_size; }; - using BuildParam = raft::neighbors::experimental::cagra::index_params; + using BuildParam = raft::neighbors::cagra::index_params; RaftCagra(Metric metric, int dim, const BuildParam& param); @@ -82,8 +82,8 @@ class RaftCagra : public ANN { private: raft::device_resources handle_; BuildParam index_params_; - raft::neighbors::experimental::cagra::search_params search_params_; - std::optional> index_; + raft::neighbors::cagra::search_params search_params_; + std::optional> index_; int device_; int dimension_; rmm::mr::pool_memory_resource mr_; @@ -105,7 +105,7 @@ template void RaftCagra::build(const T* dataset, size_t nrow, cudaStream_t) { auto dataset_view = raft::make_device_matrix_view(dataset, IdxT(nrow), dimension_); - index_.emplace(raft::neighbors::experimental::cagra::build(handle_, index_params_, dataset_view)); + index_.emplace(raft::neighbors::cagra::build(handle_, index_params_, dataset_view)); return; } @@ -118,14 +118,14 @@ void RaftCagra::set_search_param(const AnnSearchParam& param) template void RaftCagra::save(const std::string& file) const { - raft::neighbors::experimental::cagra::serialize(handle_, file, *index_); + raft::neighbors::cagra::serialize(handle_, file, *index_); return; } template void RaftCagra::load(const std::string& file) { - index_ = raft::neighbors::experimental::cagra::deserialize(handle_, file); + index_ = raft::neighbors::cagra::deserialize(handle_, file); return; } @@ -146,10 +146,10 @@ void RaftCagra::search( auto neighbors_view = raft::make_device_matrix_view(neighbors_IdxT, batch_size, k); auto distances_view = raft::make_device_matrix_view(distances, batch_size, k); - raft::neighbors::experimental::cagra::search_params search_params; + raft::neighbors::cagra::search_params search_params; search_params.max_queries = batch_size; search_params.itopk_size = search_params_.max_queries; - raft::neighbors::experimental::cagra::search( + raft::neighbors::cagra::search( handle_, search_params, *index_, queries_view, neighbors_view, distances_view); if (!std::is_same::value) { diff --git a/cpp/bench/prims/neighbors/cagra_bench.cuh b/cpp/bench/prims/neighbors/cagra_bench.cuh index 109410e399..c4a69beabb 100644 --- a/cpp/bench/prims/neighbors/cagra_bench.cuh +++ b/cpp/bench/prims/neighbors/cagra_bench.cuh @@ -73,13 +73,13 @@ struct CagraBench : public fixture { auto metric = raft::distance::DistanceType::L2Expanded; - index_.emplace(raft::neighbors::experimental::cagra::index( + index_.emplace(raft::neighbors::cagra::index( handle, metric, make_const_mdspan(dataset.view()), knn_graph.view())); } void run_benchmark(::benchmark::State& state) override { - raft::neighbors::experimental::cagra::search_params search_params; + raft::neighbors::cagra::search_params search_params; search_params.max_queries = 1024; search_params.itopk_size = params_.itopk_size; search_params.team_size = 0; @@ -95,7 +95,7 @@ struct CagraBench : public fixture { auto queries_v = make_const_mdspan(queries_.view()); loop_on_state(state, [&]() { - raft::neighbors::experimental::cagra::search( + raft::neighbors::cagra::search( this->handle, search_params, *this->index_, queries_v, ind_v, dist_v); }); @@ -123,7 +123,7 @@ struct CagraBench : public fixture { private: const params params_; - std::optional> index_; + std::optional> index_; raft::device_matrix queries_; }; diff --git a/cpp/include/raft/neighbors/cagra.cuh b/cpp/include/raft/neighbors/cagra.cuh index f093496717..a0af596d5e 100644 --- a/cpp/include/raft/neighbors/cagra.cuh +++ b/cpp/include/raft/neighbors/cagra.cuh @@ -27,7 +27,7 @@ #include #include -namespace raft::neighbors::experimental::cagra { +namespace raft::neighbors::cagra { /** * @defgroup cagra CUDA ANN Graph-based nearest neighbor search @@ -91,7 +91,7 @@ void build_knn_graph(raft::resources const& res, auto dataset_internal = mdspan, row_major, accessor>( dataset.data_handle(), dataset.extent(0), dataset.extent(1)); - detail::build_knn_graph( + cagra::detail::build_knn_graph( res, dataset_internal, knn_graph_internal, refine_rate, build_params, search_params); } @@ -149,7 +149,7 @@ void sort_knn_graph(raft::resources const& res, auto dataset_internal = mdspan, row_major, d_accessor>( dataset.data_handle(), dataset.extent(0), dataset.extent(1)); - detail::graph::sort_knn_graph(res, dataset_internal, knn_graph_internal); + cagra::detail::graph::sort_knn_graph(res, dataset_internal, knn_graph_internal); } /** @@ -188,7 +188,7 @@ void optimize(raft::resources const& res, knn_graph.extent(0), knn_graph.extent(1)); - detail::graph::optimize(res, knn_graph_internal, new_graph_internal); + cagra::detail::graph::optimize(res, knn_graph_internal, new_graph_internal); } /** @@ -312,9 +312,17 @@ void search(raft::resources const& res, auto distances_internal = raft::make_device_matrix_view( distances.data_handle(), distances.extent(0), distances.extent(1)); - detail::search_main( + cagra::detail::search_main( res, params, idx, queries_internal, neighbors_internal, distances_internal); } /** @} */ // end group cagra +} // namespace raft::neighbors::cagra + +namespace raft::neighbors::experimental::cagra { +using raft::neighbors::cagra::build; +using raft::neighbors::cagra::build_knn_graph; +using raft::neighbors::cagra::optimize; +using raft::neighbors::cagra::search; +using raft::neighbors::cagra::sort_knn_graph; } // namespace raft::neighbors::experimental::cagra diff --git a/cpp/include/raft/neighbors/cagra_serialize.cuh b/cpp/include/raft/neighbors/cagra_serialize.cuh index 8d1771a301..877cefcf98 100644 --- a/cpp/include/raft/neighbors/cagra_serialize.cuh +++ b/cpp/include/raft/neighbors/cagra_serialize.cuh @@ -18,7 +18,7 @@ #include "detail/cagra/cagra_serialize.cuh" -namespace raft::neighbors::experimental::cagra { +namespace raft::neighbors::cagra { /** * \defgroup cagra_serialize CAGRA Serialize @@ -110,7 +110,7 @@ void serialize(raft::resources const& handle, * @param[in] handle the raft handle * @param[in] is input stream * - * @return raft::neighbors::cagra::index + * @return raft::neighbors::experimental::cagra::index */ template index deserialize(raft::resources const& handle, std::istream& is) @@ -141,7 +141,7 @@ index deserialize(raft::resources const& handle, std::istream& is) * @param[in] handle the raft handle * @param[in] filename the name of the file that stores the index * - * @return raft::neighbors::cagra::index + * @return raft::neighbors::experimental::cagra::index */ template index deserialize(raft::resources const& handle, const std::string& filename) @@ -151,4 +151,10 @@ index deserialize(raft::resources const& handle, const std::string& fil /**@}*/ -} // namespace raft::neighbors::experimental::cagra +} // namespace raft::neighbors::cagra + +namespace raft::neighbors::experimental::cagra { +using raft::neighbors::cagra::deserialize; +using raft::neighbors::cagra::serialize; + +} // namespace raft::neighbors::experimental::cagra \ No newline at end of file diff --git a/cpp/include/raft/neighbors/cagra_types.hpp b/cpp/include/raft/neighbors/cagra_types.hpp index 4a384b90e1..c4e868e3a5 100644 --- a/cpp/include/raft/neighbors/cagra_types.hpp +++ b/cpp/include/raft/neighbors/cagra_types.hpp @@ -34,7 +34,7 @@ #include #include -namespace raft::neighbors::experimental::cagra { +namespace raft::neighbors::cagra { /** * @ingroup cagra * @{ @@ -229,4 +229,12 @@ struct index : ann::index { /** @} */ +} // namespace raft::neighbors::cagra + +namespace raft::neighbors::experimental::cagra { +using raft::neighbors::cagra::hash_mode; +using raft::neighbors::cagra::index; +using raft::neighbors::cagra::index_params; +using raft::neighbors::cagra::search_algo; +using raft::neighbors::cagra::search_params; } // namespace raft::neighbors::experimental::cagra diff --git a/cpp/include/raft/neighbors/detail/cagra/bitonic.hpp b/cpp/include/raft/neighbors/detail/cagra/bitonic.hpp index 45aff99421..9fca7f8ebd 100644 --- a/cpp/include/raft/neighbors/detail/cagra/bitonic.hpp +++ b/cpp/include/raft/neighbors/detail/cagra/bitonic.hpp @@ -18,7 +18,7 @@ #include #include -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace bitonic { namespace detail { @@ -223,4 +223,4 @@ __device__ void warp_sort(K k[N], V v[N], const bool asc = true) } } // namespace bitonic -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/cagra_build.cuh b/cpp/include/raft/neighbors/detail/cagra/cagra_build.cuh index 5c196471aa..421724da23 100644 --- a/cpp/include/raft/neighbors/detail/cagra/cagra_build.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/cagra_build.cuh @@ -36,7 +36,7 @@ #include #include -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { template void build_knn_graph(raft::resources const& res, @@ -229,4 +229,4 @@ void build_knn_graph(raft::resources const& res, if (!first) RAFT_LOG_DEBUG("# Finished building kNN graph"); } -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh b/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh index 1561a3bb8d..05d8b20ebc 100644 --- a/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/cagra_search.cuh @@ -30,7 +30,7 @@ #include "search_plan.cuh" #include "search_single_cta.cuh" -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { /** * @brief Search ANN using the constructed index. @@ -133,4 +133,4 @@ void search_main(raft::resources const& res, } /** @} */ // end group cagra -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/cagra_serialize.cuh b/cpp/include/raft/neighbors/detail/cagra/cagra_serialize.cuh index 7632318b88..c2526c3c18 100644 --- a/cpp/include/raft/neighbors/detail/cagra/cagra_serialize.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/cagra_serialize.cuh @@ -22,7 +22,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { // Serialization version 1. constexpr int serialization_version = 2; @@ -132,4 +132,4 @@ auto deserialize(raft::resources const& res, const std::string& filename) -> ind return index; } -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/compute_distance.hpp b/cpp/include/raft/neighbors/detail/cagra/compute_distance.hpp index f67e110fc6..91e0d88e79 100644 --- a/cpp/include/raft/neighbors/detail/cagra/compute_distance.hpp +++ b/cpp/include/raft/neighbors/detail/cagra/compute_distance.hpp @@ -22,7 +22,7 @@ #include "utils.hpp" #include -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace device { // using LOAD_256BIT_T = ulonglong4; @@ -254,4 +254,4 @@ _RAFT_DEVICE void compute_distance_to_child_nodes(INDEX_T* const result_child_in } } // namespace device -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/device_common.hpp b/cpp/include/raft/neighbors/detail/cagra/device_common.hpp index f9c81f3d25..b1a2207a4e 100644 --- a/cpp/include/raft/neighbors/detail/cagra/device_common.hpp +++ b/cpp/include/raft/neighbors/detail/cagra/device_common.hpp @@ -21,7 +21,7 @@ #include #include -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace device { // warpSize for compile time calculation @@ -49,4 +49,4 @@ _RAFT_DEVICE inline T swizzling(T x) } } // namespace device -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/factory.cuh b/cpp/include/raft/neighbors/detail/cagra/factory.cuh index 7d4cfee0b9..625040194b 100644 --- a/cpp/include/raft/neighbors/detail/cagra/factory.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/factory.cuh @@ -21,7 +21,7 @@ #include "search_plan.cuh" #include "search_single_cta.cuh" -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { template class factory { @@ -86,4 +86,4 @@ class factory { } } }; -}; // namespace raft::neighbors::experimental::cagra::detail +}; // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/fragment.hpp b/cpp/include/raft/neighbors/detail/cagra/fragment.hpp index c423ac12c2..e124b3fc8c 100644 --- a/cpp/include/raft/neighbors/detail/cagra/fragment.hpp +++ b/cpp/include/raft/neighbors/detail/cagra/fragment.hpp @@ -20,7 +20,7 @@ #include #include -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace device { namespace detail { @@ -208,4 +208,4 @@ _RAFT_DEVICE void print_fragment(const device::fragment& a) } } // namespace device -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh b/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh index d915634df9..a47e719e02 100644 --- a/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/graph_core.cuh @@ -36,7 +36,7 @@ #include "utils.hpp" -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace graph { // unnamed namespace to avoid multiple definition error @@ -588,4 +588,4 @@ void optimize(raft::resources const& res, } } // namespace graph -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/hashmap.hpp b/cpp/include/raft/neighbors/detail/cagra/hashmap.hpp index 5992aaaf1d..346bbeaa9e 100644 --- a/cpp/include/raft/neighbors/detail/cagra/hashmap.hpp +++ b/cpp/include/raft/neighbors/detail/cagra/hashmap.hpp @@ -23,7 +23,7 @@ // #pragma GCC diagnostic push // #pragma GCC diagnostic ignored // #pragma GCC diagnostic pop -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace hashmap { _RAFT_HOST_DEVICE inline uint32_t get_size(const uint32_t bitlen) { return 1U << bitlen; } @@ -85,4 +85,4 @@ _RAFT_DEVICE inline uint32_t insert(IdxT* const table, const uint32_t bitlen, co } } // namespace hashmap -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta.cuh index bf6a32eac8..24828a8c9f 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta.cuh @@ -41,7 +41,7 @@ #include #include // RAFT_CUDA_TRY_NOT_THROW is used TODO(tfeher): consider moving this to cuda_rt_essentials.hpp -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace multi_cta_search { template { }; } // namespace multi_cta_search -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh index 3ccd73d92c..4640091e69 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-ext.cuh @@ -17,7 +17,7 @@ #include // RAFT_EXPLICIT -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace multi_cta_search { #ifdef RAFT_EXPLICIT_INSTANTIATE_ONLY @@ -93,4 +93,4 @@ instantiate_kernel_selection(32, 512, uint8_t, uint32_t, float); #undef instantiate_kernel_selection } // namespace multi_cta_search -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh index 43e3e83f59..7879d61007 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh @@ -40,7 +40,7 @@ #include #include // RAFT_CUDA_TRY_NOT_THROW is used TODO(tfeher): consider moving this to cuda_rt_essentials.hpp -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace multi_cta_search { // #define _CLK_BREAKDOWN @@ -517,4 +517,4 @@ void select_and_run( // raft::resources const& res, } } // namespace multi_cta_search -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh index 033022aea1..a857d335aa 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_kernel.cuh @@ -40,7 +40,7 @@ #include #include // RAFT_CUDA_TRY_NOT_THROW is used TODO(tfeher): consider moving this to cuda_rt_essentials.hpp -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace multi_kernel_search { template @@ -738,4 +738,4 @@ struct search : search_plan_impl { }; } // namespace multi_kernel_search -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/search_plan.cuh b/cpp/include/raft/neighbors/detail/cagra/search_plan.cuh index 51ca4987d1..ae81891375 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_plan.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_plan.cuh @@ -26,7 +26,7 @@ #include #include -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { struct search_plan_impl_base : public search_params { int64_t max_dim; @@ -325,4 +325,4 @@ struct search_plan_impl : public search_plan_impl_base { // }; /** @} */ // end group cagra -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/search_single_cta.cuh b/cpp/include/raft/neighbors/detail/cagra/search_single_cta.cuh index bad2039f8c..9fc97facda 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_single_cta.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_single_cta.cuh @@ -42,7 +42,7 @@ #include #include // RAFT_CUDA_TRY_NOT_THROW is used TODO(tfeher): consider moving this to cuda_rt_essentials.hpp -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace single_cta_search { template { }; } // namespace single_cta_search -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh index b0130e45d4..f589fd4637 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-ext.cuh @@ -16,7 +16,7 @@ #pragma once #include // RAFT_EXPLICIT -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace single_cta_search { #ifdef RAFT_EXPLICIT_INSTANTIATE_ONLY @@ -98,4 +98,4 @@ instantiate_single_cta_select_and_run(32, 512, uint8_t, uint32_t, float); #undef instantiate_single_cta_select_and_run } // namespace single_cta_search -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh index ca2166ab8d..df822c0113 100644 --- a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh @@ -41,7 +41,7 @@ #include #include // RAFT_CUDA_TRY_NOT_THROW is used TODO(tfeher): consider moving this to cuda_rt_essentials.hpp -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace single_cta_search { // #define _CLK_BREAKDOWN @@ -887,4 +887,4 @@ void select_and_run( // raft::resources const& res, RAFT_CUDA_TRY(cudaPeekAtLastError()); } } // namespace single_cta_search -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/topk_by_radix.cuh b/cpp/include/raft/neighbors/detail/cagra/topk_by_radix.cuh index d151cc8ee7..a1b7f930d3 100644 --- a/cpp/include/raft/neighbors/detail/cagra/topk_by_radix.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/topk_by_radix.cuh @@ -17,7 +17,7 @@ #include "topk_for_cagra/topk_core.cuh" -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace single_cta_search { template @@ -94,4 +94,4 @@ TOP_FUNC_PARTIAL_SPECIALIZATION(512); TOP_FUNC_PARTIAL_SPECIALIZATION(1024); } // namespace single_cta_search -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/topk_for_cagra/topk.h b/cpp/include/raft/neighbors/detail/cagra/topk_for_cagra/topk.h index 2896dba1f3..92b9474047 100644 --- a/cpp/include/raft/neighbors/detail/cagra/topk_for_cagra/topk.h +++ b/cpp/include/raft/neighbors/detail/cagra/topk_for_cagra/topk.h @@ -18,7 +18,7 @@ #include #include -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { // size_t _cuann_find_topk_bufferSize(uint32_t topK, @@ -55,4 +55,4 @@ CUDA_DEVICE_HOST_FUNC inline size_t _cuann_aligned(size_t size, size_t unit = 12 if (size % unit) { size += unit - (size % unit); } return size; } -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/topk_for_cagra/topk_core.cuh b/cpp/include/raft/neighbors/detail/cagra/topk_for_cagra/topk_core.cuh index 9faf57c0f5..dd73558f86 100644 --- a/cpp/include/raft/neighbors/detail/cagra/topk_for_cagra/topk_core.cuh +++ b/cpp/include/raft/neighbors/detail/cagra/topk_for_cagra/topk_core.cuh @@ -21,7 +21,7 @@ #include #include -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { using namespace cub; // @@ -927,4 +927,4 @@ inline void _cuann_find_topk(uint32_t topK, return; } -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/include/raft/neighbors/detail/cagra/utils.hpp b/cpp/include/raft/neighbors/detail/cagra/utils.hpp index 934e84d4d5..22c7a60647 100644 --- a/cpp/include/raft/neighbors/detail/cagra/utils.hpp +++ b/cpp/include/raft/neighbors/detail/cagra/utils.hpp @@ -22,7 +22,7 @@ #include #include -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace utils { template inline cudaDataType_t get_cuda_data_type(); @@ -150,4 +150,4 @@ struct gen_index_msb_1_mask { }; } // namespace utils -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_00_generate.py b/cpp/src/neighbors/detail/cagra/search_multi_cta_00_generate.py index 170c57c521..32bec82d38 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 @@ -40,7 +40,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \\ template void select_and_run( \\ @@ -73,7 +73,7 @@ trailer = """ #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::namespace multi_cta_search +} // namespace raft::neighbors::cagra::detail::namespace multi_cta_search """ mxdim_team = [(128, 8), (256, 16), (512, 32), (1024, 32)] diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim1024_t32.cu index 207028dcec..7536c8e9d5 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(32, 1024, float, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim128_t8.cu index 4a5c0f106b..96b4dab650 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(8, 128, float, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim256_t16.cu index 93a9f41881..410ac55f66 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(16, 256, float, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint32_dim512_t32.cu index fb321b2cf7..f80b8603d1 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(32, 512, float, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim1024_t32.cu index e73698460d..97f29da4c2 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(32, 1024, float, uint64_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim128_t8.cu index e51fdcbc62..959e36ed7e 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(8, 128, float, uint64_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim256_t16.cu index caa45b5395..4324df905b 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(16, 256, float, uint64_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_float_uint64_dim512_t32.cu index 67e54f0937..e1d1f8fa62 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(32, 512, float, uint64_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim1024_t32.cu index 2e929eb4f0..08fbdbcf5c 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(32, 1024, int8_t, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim128_t8.cu index d3e2e78250..e4015dfbd3 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(8, 128, int8_t, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim256_t16.cu index 802edafdf2..22622b380a 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(16, 256, int8_t, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_int8_uint32_dim512_t32.cu index 96e91c475e..8fcd2008c0 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(32, 512, int8_t, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim1024_t32.cu index 6db346c67a..feb9b01819 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(32, 1024, uint8_t, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim128_t8.cu index 4b1c6c89f4..7fa2d4f1a8 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(8, 128, uint8_t, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim256_t16.cu index f978a9011a..8f278f9b4b 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(16, 256, uint8_t, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_multi_cta_uint8_uint32_dim512_t32.cu index 390330ec93..cec0753442 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::multi_cta_search { +namespace raft::neighbors::cagra::detail::multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ template void select_and_run( \ @@ -58,4 +58,4 @@ instantiate_kernel_selection(32, 512, uint8_t, uint32_t, float); #undef instantiate_kernel_selection -} // namespace raft::neighbors::experimental::cagra::detail::multi_cta_search +} // namespace raft::neighbors::cagra::detail::multi_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_00_generate.py b/cpp/src/neighbors/detail/cagra/search_single_cta_00_generate.py index b8f623d4c4..ba6ce82485 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 @@ -40,7 +40,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \\ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \\ @@ -75,7 +75,7 @@ trailer = """ #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search """ mxdim_team = [(128, 8), (256, 16), (512, 32), (1024, 32)] diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim1024_t32.cu index 523f2761fc..7474875bf9 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(32, 1024, float, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim128_t8.cu index cb8b21bfe8..8efd7ade82 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(8, 128, float, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim256_t16.cu index f5ccfa7572..df88617904 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(16, 256, float, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint32_dim512_t32.cu index 1d83979a88..fb9bf8b7d5 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(32, 512, float, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu index cd588e13ef..da49fa2f4b 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(32, 1024, float, uint64_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim128_t8.cu index b47db68273..3c5e595329 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(8, 128, float, uint64_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim256_t16.cu index d875080345..a32d2f4516 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(16, 256, float, uint64_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_float_uint64_dim512_t32.cu index 848e71a645..1efcbcc125 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(32, 512, float, uint64_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim1024_t32.cu index de7acb56fe..8e3b2ed6f6 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(32, 1024, int8_t, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim128_t8.cu index d0e90603e2..ad3db811ec 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(8, 128, int8_t, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim256_t16.cu index 26764c5ad9..845ee65c33 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(16, 256, int8_t, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_int8_uint32_dim512_t32.cu index 6568ab6dba..c0a237140c 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(32, 512, int8_t, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim1024_t32.cu index 311f42c9a7..07e678bcb5 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(32, 1024, uint8_t, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim128_t8.cu index 197aa71d7b..33a956e77d 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(8, 128, uint8_t, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim256_t16.cu index dfb47a1137..cfc4598404 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(16, 256, uint8_t, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu b/cpp/src/neighbors/detail/cagra/search_single_cta_uint8_uint32_dim512_t32.cu index 1b874bcf9b..ee4897ff3f 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 @@ -26,7 +26,7 @@ #include -namespace raft::neighbors::experimental::cagra::detail::single_cta_search { +namespace raft::neighbors::cagra::detail::single_cta_search { #define instantiate_single_cta_select_and_run( \ TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -60,4 +60,4 @@ instantiate_single_cta_select_and_run(32, 512, uint8_t, uint32_t, float); #undef instantiate_single_cta_search_kernel -} // namespace raft::neighbors::experimental::cagra::detail::single_cta_search +} // namespace raft::neighbors::cagra::detail::single_cta_search diff --git a/cpp/test/neighbors/ann_cagra.cuh b/cpp/test/neighbors/ann_cagra.cuh index 3e929f9f3b..77e355981f 100644 --- a/cpp/test/neighbors/ann_cagra.cuh +++ b/cpp/test/neighbors/ann_cagra.cuh @@ -52,9 +52,9 @@ void RandomSuffle(raft::host_matrix_view index) IdxT* const row_ptr = index.data_handle() + i * index.extent(1); for (unsigned j = 0; j < index.extent(1); j++) { // Swap two indices at random - rand = raft::neighbors::experimental::cagra::detail::device::xorshift64(rand); + rand = raft::neighbors::cagra::detail::device::xorshift64(rand); const auto i0 = rand % index.extent(1); - rand = raft::neighbors::experimental::cagra::detail::device::xorshift64(rand); + rand = raft::neighbors::cagra::detail::device::xorshift64(rand); const auto i1 = rand % index.extent(1); const auto tmp = row_ptr[i0]; diff --git a/cpp/test/neighbors/ann_cagra/search_kernel_uint64_t.cuh b/cpp/test/neighbors/ann_cagra/search_kernel_uint64_t.cuh index 562e5ac2ca..53278b9666 100644 --- a/cpp/test/neighbors/ann_cagra/search_kernel_uint64_t.cuh +++ b/cpp/test/neighbors/ann_cagra/search_kernel_uint64_t.cuh @@ -17,7 +17,7 @@ #include // RAFT_EXPLICIT -namespace raft::neighbors::experimental::cagra::detail { +namespace raft::neighbors::cagra::detail { namespace multi_cta_search { #define instantiate_kernel_selection(TEAM_SIZE, MAX_DATASET_DIM, DATA_T, INDEX_T, DISTANCE_T) \ @@ -90,4 +90,4 @@ instantiate_single_cta_select_and_run(16, 256, float, uint64_t, float); instantiate_single_cta_select_and_run(32, 512, float, uint64_t, float); } // namespace single_cta_search -} // namespace raft::neighbors::experimental::cagra::detail +} // namespace raft::neighbors::cagra::detail diff --git a/docs/source/cpp_api/neighbors_cagra.rst b/docs/source/cpp_api/neighbors_cagra.rst index 68372bbb71..6613b0b06d 100644 --- a/docs/source/cpp_api/neighbors_cagra.rst +++ b/docs/source/cpp_api/neighbors_cagra.rst @@ -11,7 +11,7 @@ Please note that the CAGRA implementation is currently experimental and the API ``#include `` -namespace *raft::neighbors::experimental::cagra* +namespace *raft::neighbors::cagra* .. doxygengroup:: cagra :project: RAFT From 410f0b9a4502b368747694a30b9bc48a42817500 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Mon, 24 Jul 2023 13:25:51 -0400 Subject: [PATCH 2/3] Adding TODOs in cagra public APIs to remove experimental namespaces --- cpp/include/raft/neighbors/cagra.cuh | 1 + cpp/include/raft/neighbors/cagra_serialize.cuh | 1 + cpp/include/raft/neighbors/cagra_types.hpp | 1 + 3 files changed, 3 insertions(+) diff --git a/cpp/include/raft/neighbors/cagra.cuh b/cpp/include/raft/neighbors/cagra.cuh index a0af596d5e..5f9c8ceaa8 100644 --- a/cpp/include/raft/neighbors/cagra.cuh +++ b/cpp/include/raft/neighbors/cagra.cuh @@ -319,6 +319,7 @@ void search(raft::resources const& res, } // namespace raft::neighbors::cagra +// TODO: Remove deprecated experimental namespace in 23.12 release namespace raft::neighbors::experimental::cagra { using raft::neighbors::cagra::build; using raft::neighbors::cagra::build_knn_graph; diff --git a/cpp/include/raft/neighbors/cagra_serialize.cuh b/cpp/include/raft/neighbors/cagra_serialize.cuh index 877cefcf98..2242629409 100644 --- a/cpp/include/raft/neighbors/cagra_serialize.cuh +++ b/cpp/include/raft/neighbors/cagra_serialize.cuh @@ -153,6 +153,7 @@ index deserialize(raft::resources const& handle, const std::string& fil } // namespace raft::neighbors::cagra +// TODO: Remove deprecated experimental namespace in 23.12 release namespace raft::neighbors::experimental::cagra { using raft::neighbors::cagra::deserialize; using raft::neighbors::cagra::serialize; diff --git a/cpp/include/raft/neighbors/cagra_types.hpp b/cpp/include/raft/neighbors/cagra_types.hpp index c4e868e3a5..b943418fc1 100644 --- a/cpp/include/raft/neighbors/cagra_types.hpp +++ b/cpp/include/raft/neighbors/cagra_types.hpp @@ -231,6 +231,7 @@ struct index : ann::index { } // namespace raft::neighbors::cagra +// TODO: Remove deprecated experimental namespace in 23.12 release namespace raft::neighbors::experimental::cagra { using raft::neighbors::cagra::hash_mode; using raft::neighbors::cagra::index; From 987a09adda36f9bdd1f1a8be885cd3395ea2c1d5 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Mon, 24 Jul 2023 13:36:29 -0400 Subject: [PATCH 3/3] Fixing style for conflicts resolved through github --- cpp/bench/ann/src/raft/raft_cagra_wrapper.h | 7 +--- cpp/bench/prims/neighbors/cagra_bench.cuh | 1 - cpp/test/CMakeLists.txt | 44 +++++++-------------- python/pylibraft/CMakeLists.txt | 4 +- 4 files changed, 18 insertions(+), 38 deletions(-) diff --git a/cpp/bench/ann/src/raft/raft_cagra_wrapper.h b/cpp/bench/ann/src/raft/raft_cagra_wrapper.h index 9f4483f01d..e898a13636 100644 --- a/cpp/bench/ann/src/raft/raft_cagra_wrapper.h +++ b/cpp/bench/ann/src/raft/raft_cagra_wrapper.h @@ -106,13 +106,11 @@ void RaftCagra::build(const T* dataset, size_t nrow, cudaStream_t) { if (get_property().dataset_memory_type == MemoryType::Host) { auto dataset_view = raft::make_host_matrix_view(dataset, IdxT(nrow), dimension_); - index_.emplace( - raft::neighbors::cagra::build(handle_, index_params_, dataset_view)); + index_.emplace(raft::neighbors::cagra::build(handle_, index_params_, dataset_view)); } else { auto dataset_view = raft::make_device_matrix_view(dataset, IdxT(nrow), dimension_); - index_.emplace( - raft::neighbors::cagra::build(handle_, index_params_, dataset_view)); + index_.emplace(raft::neighbors::cagra::build(handle_, index_params_, dataset_view)); } return; } @@ -156,7 +154,6 @@ void RaftCagra::search( auto neighbors_view = raft::make_device_matrix_view(neighbors_IdxT, batch_size, k); auto distances_view = raft::make_device_matrix_view(distances, batch_size, k); - raft::neighbors::cagra::search( handle_, search_params_, *index_, queries_view, neighbors_view, distances_view); diff --git a/cpp/bench/prims/neighbors/cagra_bench.cuh b/cpp/bench/prims/neighbors/cagra_bench.cuh index ee096dfff2..dbc6645f92 100644 --- a/cpp/bench/prims/neighbors/cagra_bench.cuh +++ b/cpp/bench/prims/neighbors/cagra_bench.cuh @@ -74,7 +74,6 @@ struct CagraBench : public fixture { auto metric = raft::distance::DistanceType::L2Expanded; - index_.emplace(raft::neighbors::cagra::index( handle, metric, make_const_mdspan(dataset_.view()), make_const_mdspan(knn_graph_.view()))); } diff --git a/cpp/test/CMakeLists.txt b/cpp/test/CMakeLists.txt index 87fc7ed782..41dab0e388 100644 --- a/cpp/test/CMakeLists.txt +++ b/cpp/test/CMakeLists.txt @@ -54,8 +54,7 @@ function(ConfigureTest) ) set_target_properties( ${TEST_NAME} - PROPERTIES - RUNTIME_OUTPUT_DIRECTORY "$" + PROPERTIES RUNTIME_OUTPUT_DIRECTORY "$" INSTALL_RPATH "\$ORIGIN/../../../lib" CXX_STANDARD 17 CXX_STANDARD_REQUIRED ON @@ -189,9 +188,7 @@ if(BUILD_TESTS) NAME EXT_HEADERS_TEST_COMPILED_EXPLICIT PATH ${EXT_HEADER_TEST_SOURCES} LIB EXPLICIT_INSTANTIATE_ONLY ) - ConfigureTest( - NAME EXT_HEADERS_TEST_COMPILED_IMPLICIT PATH ${EXT_HEADER_TEST_SOURCES} LIB - ) + ConfigureTest(NAME EXT_HEADERS_TEST_COMPILED_IMPLICIT PATH ${EXT_HEADER_TEST_SOURCES} LIB) ConfigureTest(NAME EXT_HEADERS_TEST_IMPLICIT PATH ${EXT_HEADER_TEST_SOURCES}) ConfigureTest(NAME LABEL_TEST PATH test/label/label.cu test/label/merge_labels.cu) @@ -256,26 +253,12 @@ if(BUILD_TESTS) EXPLICIT_INSTANTIATE_ONLY ) + ConfigureTest(NAME MATRIX_SELECT_TEST PATH test/matrix/select_k.cu LIB EXPLICIT_INSTANTIATE_ONLY) ConfigureTest( - NAME - MATRIX_SELECT_TEST - PATH - test/matrix/select_k.cu - LIB - EXPLICIT_INSTANTIATE_ONLY - ) - - ConfigureTest( - NAME - MATRIX_SELECT_LARGE_TEST - PATH - test/matrix/select_large_k.cu - LIB - EXPLICIT_INSTANTIATE_ONLY + NAME MATRIX_SELECT_LARGE_TEST PATH test/matrix/select_large_k.cu LIB EXPLICIT_INSTANTIATE_ONLY ) - ConfigureTest( NAME RANDOM_TEST @@ -319,7 +302,7 @@ if(BUILD_TESTS) ConfigureTest( NAME SPARSE_DIST_TEST PATH test/sparse/dist_coo_spmv.cu test/sparse/distance.cu - test/sparse/gram.cu LIB EXPLICIT_INSTANTIATE_ONLY + test/sparse/gram.cu LIB EXPLICIT_INSTANTIATE_ONLY ) ConfigureTest( @@ -366,7 +349,10 @@ if(BUILD_TESTS) src/neighbors/detail/cagra/search_single_cta_float_uint64_dim1024_t32.cu LIB EXPLICIT_INSTANTIATE_ONLY - GPUS 1 PERCENT 100 + GPUS + 1 + PERCENT + 100 ) ConfigureTest( @@ -383,16 +369,14 @@ if(BUILD_TESTS) test/neighbors/ann_ivf_pq/test_uint8_t_int64_t.cu LIB EXPLICIT_INSTANTIATE_ONLY - GPUS 1 PERCENT 100 + GPUS + 1 + PERCENT + 100 ) ConfigureTest( - NAME - NEIGHBORS_SELECTION_TEST - PATH - test/neighbors/selection.cu - LIB - EXPLICIT_INSTANTIATE_ONLY + NAME NEIGHBORS_SELECTION_TEST PATH test/neighbors/selection.cu LIB EXPLICIT_INSTANTIATE_ONLY GPUS 1 PERCENT 50 ) diff --git a/python/pylibraft/CMakeLists.txt b/python/pylibraft/CMakeLists.txt index 91e52e962f..29405e43c0 100644 --- a/python/pylibraft/CMakeLists.txt +++ b/python/pylibraft/CMakeLists.txt @@ -18,8 +18,8 @@ include(../../fetch_rapids.cmake) set(pylibraft_version 23.08.00) -# We always need CUDA for pylibraft because the raft dependency brings in a -# header-only cuco dependency that enables CUDA unconditionally. +# We always need CUDA for pylibraft because the raft dependency brings in a header-only cuco +# dependency that enables CUDA unconditionally. include(rapids-cuda) rapids_cuda_init_architectures(pylibraft)