diff --git a/c/src/neighbors/brute_force.cpp b/c/src/neighbors/brute_force.cpp index 081f433a3e..5a6a7500f5 100644 --- a/c/src/neighbors/brute_force.cpp +++ b/c/src/neighbors/brute_force.cpp @@ -1,6 +1,6 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -250,10 +250,13 @@ extern "C" cuvsError_t cuvsBruteForceDeserialize(cuvsResources_t res, index->dtype.bits = dtype.itemsize * 8; if (dtype.kind == 'f' && dtype.itemsize == 4) { index->dtype.code = kDLFloat; - index->addr = reinterpret_cast(_deserialize(res, filename)); - } else if (dtype.kind == 'e' && dtype.itemsize == 2) { + index->addr = + reinterpret_cast(_deserialize(res, filename)); + } else if ((dtype.kind == 'f' || dtype.kind == 'e') && + dtype.itemsize == 2) { index->dtype.code = kDLFloat; - index->addr = reinterpret_cast(_deserialize(res, filename)); + index->addr = + reinterpret_cast(_deserialize(res, filename)); } else { RAFT_FAIL("Unsupported index dtype: %d and bits: %d", index->dtype.code, index->dtype.bits); } diff --git a/c/src/neighbors/cagra.cpp b/c/src/neighbors/cagra.cpp index 776b4b07b0..57e8168c1e 100644 --- a/c/src/neighbors/cagra.cpp +++ b/c/src/neighbors/cagra.cpp @@ -1056,7 +1056,7 @@ static auto read_serialized_header(cuvsResources_t res, const char *filename) .code = 0, .bits = static_cast(dtype.itemsize * 8), .lanes = 1}; if (dtype.kind == 'f' && dtype.itemsize == 4) { output_dtype.code = kDLFloat; - } else if (dtype.kind == 'e' && dtype.itemsize == 2) { + } else if ((dtype.kind == 'f' || dtype.kind == 'e') && dtype.itemsize == 2) { output_dtype.code = kDLFloat; } else if (dtype.kind == 'i' && dtype.itemsize == 1) { output_dtype.code = kDLInt; diff --git a/c/src/neighbors/ivf_flat.cpp b/c/src/neighbors/ivf_flat.cpp index 729c810aa0..bd0e67869e 100644 --- a/c/src/neighbors/ivf_flat.cpp +++ b/c/src/neighbors/ivf_flat.cpp @@ -1,6 +1,6 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -310,10 +310,13 @@ extern "C" cuvsError_t cuvsIvfFlatDeserialize(cuvsResources_t res, index->dtype.bits = dtype.itemsize * 8; if (dtype.kind == 'f' && dtype.itemsize == 4) { - index->addr = reinterpret_cast(_deserialize(res, filename)); + index->addr = reinterpret_cast( + _deserialize(res, filename)); index->dtype.code = kDLFloat; - } else if (dtype.kind == 'e' && dtype.itemsize == 2) { - index->addr = reinterpret_cast(_deserialize(res, filename)); + } else if ((dtype.kind == 'f' || dtype.kind == 'e') && + dtype.itemsize == 2) { + index->addr = reinterpret_cast( + _deserialize(res, filename)); index->dtype.code = kDLFloat; index->dtype.bits = 16; } else if (dtype.kind == 'i' && dtype.itemsize == 1) { diff --git a/c/src/neighbors/mg_cagra.cpp b/c/src/neighbors/mg_cagra.cpp index 79caf18593..664883cd6f 100644 --- a/c/src/neighbors/mg_cagra.cpp +++ b/c/src/neighbors/mg_cagra.cpp @@ -581,7 +581,7 @@ extern "C" cuvsError_t cuvsMultiGpuCagraDeserialize(cuvsResources_t res, if (dtype.kind == 'f' && dtype.itemsize == 4) { index->dtype.code = kDLFloat; index->addr = try_layout_deser(float{}); - } else if (dtype.kind == 'e' && dtype.itemsize == 2) { + } else if ((dtype.kind == 'f' || dtype.kind == 'e') && dtype.itemsize == 2) { index->dtype.code = kDLFloat; index->addr = try_layout_deser(half{}); } else if (dtype.kind == 'i' && dtype.itemsize == 1) { @@ -627,7 +627,7 @@ extern "C" cuvsError_t cuvsMultiGpuCagraDistribute(cuvsResources_t res, if (dtype.kind == 'f' && dtype.itemsize == 4) { index->dtype.code = kDLFloat; index->addr = try_layout_distribute(float{}); - } else if (dtype.kind == 'e' && dtype.itemsize == 2) { + } else if ((dtype.kind == 'f' || dtype.kind == 'e') && dtype.itemsize == 2) { index->dtype.code = kDLFloat; index->addr = try_layout_distribute(half{}); } else if (dtype.kind == 'i' && dtype.itemsize == 1) { diff --git a/c/src/neighbors/mg_ivf_flat.cpp b/c/src/neighbors/mg_ivf_flat.cpp index 4e1b2883ea..410ec66185 100644 --- a/c/src/neighbors/mg_ivf_flat.cpp +++ b/c/src/neighbors/mg_ivf_flat.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -408,10 +408,13 @@ extern "C" cuvsError_t cuvsMultiGpuIvfFlatDeserialize(cuvsResources_t res, index->dtype.bits = dtype.itemsize * 8; if (dtype.kind == 'f' && dtype.itemsize == 4) { index->dtype.code = kDLFloat; - index->addr = reinterpret_cast(_mg_deserialize(res, filename)); - } else if (dtype.kind == 'e' && dtype.itemsize == 2) { + index->addr = + reinterpret_cast(_mg_deserialize(res, filename)); + } else if ((dtype.kind == 'f' || dtype.kind == 'e') && + dtype.itemsize == 2) { index->dtype.code = kDLFloat; - index->addr = reinterpret_cast(_mg_deserialize(res, filename)); + index->addr = + reinterpret_cast(_mg_deserialize(res, filename)); } else if (dtype.kind == 'i' && dtype.itemsize == 1) { index->dtype.code = kDLInt; index->addr = reinterpret_cast(_mg_deserialize(res, filename)); @@ -442,10 +445,13 @@ extern "C" cuvsError_t cuvsMultiGpuIvfFlatDistribute(cuvsResources_t res, index->dtype.bits = dtype.itemsize * 8; if (dtype.kind == 'f' && dtype.itemsize == 4) { index->dtype.code = kDLFloat; - index->addr = reinterpret_cast(_mg_distribute(res, filename)); - } else if (dtype.kind == 'e' && dtype.itemsize == 2) { + index->addr = + reinterpret_cast(_mg_distribute(res, filename)); + } else if ((dtype.kind == 'f' || dtype.kind == 'e') && + dtype.itemsize == 2) { index->dtype.code = kDLFloat; - index->addr = reinterpret_cast(_mg_distribute(res, filename)); + index->addr = + reinterpret_cast(_mg_distribute(res, filename)); } else if (dtype.kind == 'i' && dtype.itemsize == 1) { index->dtype.code = kDLInt; index->addr = reinterpret_cast(_mg_distribute(res, filename)); diff --git a/c/src/neighbors/mg_ivf_pq.cpp b/c/src/neighbors/mg_ivf_pq.cpp index 41aa323138..56c2aabc65 100644 --- a/c/src/neighbors/mg_ivf_pq.cpp +++ b/c/src/neighbors/mg_ivf_pq.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -400,10 +400,13 @@ extern "C" cuvsError_t cuvsMultiGpuIvfPqDeserialize(cuvsResources_t res, index->dtype.bits = dtype.itemsize * 8; if (dtype.kind == 'f' && dtype.itemsize == 4) { index->dtype.code = kDLFloat; - index->addr = reinterpret_cast(_mg_deserialize(res, filename)); - } else if (dtype.kind == 'e' && dtype.itemsize == 2) { + index->addr = + reinterpret_cast(_mg_deserialize(res, filename)); + } else if ((dtype.kind == 'f' || dtype.kind == 'e') && + dtype.itemsize == 2) { index->dtype.code = kDLFloat; - index->addr = reinterpret_cast(_mg_deserialize(res, filename)); + index->addr = + reinterpret_cast(_mg_deserialize(res, filename)); } else if (dtype.kind == 'i' && dtype.itemsize == 1) { index->dtype.code = kDLInt; index->addr = reinterpret_cast(_mg_deserialize(res, filename)); diff --git a/cpp/bench/ann/src/cuvs/cuvs_cagra_hnswlib.cu b/cpp/bench/ann/src/cuvs/cuvs_cagra_hnswlib.cu index c903b39fcc..7636e39a56 100644 --- a/cpp/bench/ann/src/cuvs/cuvs_cagra_hnswlib.cu +++ b/cpp/bench/ann/src/cuvs/cuvs_cagra_hnswlib.cu @@ -1,9 +1,10 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "../common/ann_types.hpp" +#include "../common/conf.hpp" #include "cuvs_ann_bench_param_parser.h" #include "cuvs_cagra_hnswlib_wrapper.h" @@ -33,15 +34,71 @@ auto parse_build_param(const nlohmann::json& conf) -> } else { hnsw_params.hierarchy = cuvs::neighbors::hnsw::HnswHierarchy::GPU; } + if (conf.contains("output_format")) { + if (conf.at("output_format") == "hnswlib") { + hnsw_params.output_format = cuvs::neighbors::hnsw::HnswOutputFormat::HNSWLIB; + } else if (conf.at("output_format") == "graph_only") { + hnsw_params.output_format = cuvs::neighbors::hnsw::HnswOutputFormat::GRAPH_ONLY; + } else { + THROW("Invalid value for output_format: %s", + conf.at("output_format").get().c_str()); + } + } if (conf.contains("ef_construction")) { hnsw_params.ef_construction = conf.at("ef_construction"); } + if (conf.contains("dataset_path")) { + param.dataset_path = conf.at("dataset_path"); + } else if (hnsw_params.output_format == cuvs::neighbors::hnsw::HnswOutputFormat::GRAPH_ONLY) { + param.dataset_path = configuration::singleton().get_dataset_conf().base_file; + } if (conf.contains("num_threads")) { hnsw_params.num_threads = conf.at("num_threads"); } // Reuse the CAGRA wrapper params parser ::parse_build_param(conf, cagra_params); - if (conf.contains("M")) { hnsw_params.M = conf.at("M"); } + + // ACE / GRAPH_ONLY builds can be fine-tuned from the benchmark config. The library + // auto-selects the build algorithm from `M` and `ef_construction`; here we only forward the + // explicit ACE overrides (if any) onto the new hnsw index params. + auto ace_conf = collect_conf_with_prefix(conf, "ace_"); + if (!ace_conf.empty()) { + auto ace_params = cuvs::neighbors::hnsw::graph_build_params::ace_params(); + if (ace_conf.contains("npartitions")) { ace_params.npartitions = ace_conf.at("npartitions"); } + if (ace_conf.contains("build_dir")) { ace_params.build_dir = ace_conf.at("build_dir"); } + if (ace_conf.contains("ef_construction")) { + ace_params.ef_construction = ace_conf.at("ef_construction"); + } + if (ace_conf.contains("use_disk")) { ace_params.use_disk = ace_conf.at("use_disk"); } + hnsw_params.graph_build_params = ace_params; + } + + // GRAPH_ONLY always needs disk-backed ACE settings before hnsw::build. + if (hnsw_params.output_format == cuvs::neighbors::hnsw::HnswOutputFormat::GRAPH_ONLY) { + auto ace_params = std::holds_alternative( + hnsw_params.graph_build_params) + ? std::get( + hnsw_params.graph_build_params) + : cuvs::neighbors::hnsw::graph_build_params::ace_params(); + if (!ace_conf.contains("use_disk")) { ace_params.use_disk = true; } + const auto use_disk_conf = + ace_conf.contains("use_disk") ? ace_conf.at("use_disk").dump() : std::string{"unset"}; + const auto build_dir_conf = + ace_conf.contains("build_dir") ? ace_conf.at("build_dir").dump() : std::string{"unset"}; + RAFT_EXPECTS(ace_params.use_disk, + "GRAPH_ONLY requires ACE disk mode (ace_params.use_disk = true); " + "got ace_use_disk=%s", + use_disk_conf.c_str()); + RAFT_EXPECTS(!ace_params.build_dir.empty(), + "GRAPH_ONLY requires ace_params.build_dir to be set; " + "got ace_build_dir=%s", + build_dir_conf.c_str()); + RAFT_EXPECTS(!param.dataset_path.empty(), + "GRAPH_ONLY requires dataset_path or a configured dataset base_file; " + "got dataset_path='%s'", + param.dataset_path.c_str()); + hnsw_params.graph_build_params = ace_params; + } return param; } diff --git a/cpp/bench/ann/src/cuvs/cuvs_cagra_hnswlib_wrapper.h b/cpp/bench/ann/src/cuvs/cuvs_cagra_hnswlib_wrapper.h index 81c0e60ecf..b2cec9529e 100644 --- a/cpp/bench/ann/src/cuvs/cuvs_cagra_hnswlib_wrapper.h +++ b/cpp/bench/ann/src/cuvs/cuvs_cagra_hnswlib_wrapper.h @@ -9,10 +9,34 @@ #include #include +#include #include +#include namespace cuvs::bench { +inline void move_file_overwrite(const std::filesystem::path& src, const std::filesystem::path& dst) +{ + std::error_code ec; + if (src == dst || + (std::filesystem::exists(dst, ec) && std::filesystem::equivalent(src, dst, ec))) { + return; + } + if (!dst.parent_path().empty()) { std::filesystem::create_directories(dst.parent_path()); } + if (std::filesystem::exists(dst, ec)) { std::filesystem::remove(dst, ec); } + + std::filesystem::rename(src, dst, ec); + if (ec) { + // Rename fails across filesystems. Fall back to copy followed by removal of the source. + ec.clear(); + std::filesystem::copy_file(src, dst, std::filesystem::copy_options::overwrite_existing, ec); + const auto src_str = src.string(); + const auto dst_str = dst.string(); + RAFT_EXPECTS(!ec, "Failed to move '%s' to '%s'.", src_str.c_str(), dst_str.c_str()); + std::filesystem::remove(src, ec); + } +} + template class cuvs_cagra_hnswlib : public algo, public algo_gpu { public: @@ -22,6 +46,7 @@ class cuvs_cagra_hnswlib : public algo, public algo_gpu { using cagra_wrapper_params = typename cuvs_cagra::build_param; cagra_wrapper_params cagra_build_params; cuvs::neighbors::hnsw::index_params hnsw_index_params; + std::string dataset_path; }; struct search_param : public search_param_base { @@ -97,18 +122,25 @@ void cuvs_cagra_hnswlib::set_search_param(const search_param_base& para template void cuvs_cagra_hnswlib::save(const std::string& file) const { + if (build_param_.hnsw_index_params.output_format == + cuvs::neighbors::hnsw::HnswOutputFormat::GRAPH_ONLY) { + const auto src_artifact = std::filesystem::path(hnsw_index_->file_path()); + RAFT_EXPECTS(!src_artifact.empty(), "Layered HNSW artifact path is not available."); + RAFT_EXPECTS(std::filesystem::exists(src_artifact), + "Layered HNSW artifact '%s' does not exist.", + src_artifact.c_str()); + + move_file_overwrite(src_artifact, std::filesystem::path(file)); + return; + } + if (cagra_ace_build_) { std::string index_filename = hnsw_index_->file_path(); RAFT_EXPECTS(!index_filename.empty(), "HNSW index file path is not available."); RAFT_EXPECTS(std::filesystem::exists(index_filename), "Index file '%s' does not exist.", index_filename.c_str()); - if (std::filesystem::exists(file)) { std::filesystem::remove(file); } - // might fail when using 2 different filesystems - std::error_code ec; - std::filesystem::rename(index_filename, file, ec); - RAFT_EXPECTS( - !ec, "Failed to rename index file '%s' to '%s'.", index_filename.c_str(), file.c_str()); + move_file_overwrite(std::filesystem::path(index_filename), std::filesystem::path(file)); } else { cuvs::neighbors::hnsw::serialize(handle_, file, *(hnsw_index_.get())); } @@ -118,12 +150,17 @@ template void cuvs_cagra_hnswlib::load(const std::string& file) { cuvs::neighbors::hnsw::index* idx = nullptr; - cuvs::neighbors::hnsw::deserialize(handle_, - build_param_.hnsw_index_params, - file, - this->dim_, - parse_metric_type(this->metric_), - &idx); + if (build_param_.hnsw_index_params.output_format == + cuvs::neighbors::hnsw::HnswOutputFormat::GRAPH_ONLY) { + cuvs::neighbors::hnsw::deserialize(handle_, file, build_param_.dataset_path, &idx); + } else { + cuvs::neighbors::hnsw::deserialize(handle_, + build_param_.hnsw_index_params, + file, + this->dim_, + parse_metric_type(this->metric_), + &idx); + } hnsw_index_ = std::shared_ptr>(idx); } diff --git a/cpp/include/cuvs/neighbors/hnsw.hpp b/cpp/include/cuvs/neighbors/hnsw.hpp index d5215637c2..fe412e8cc1 100644 --- a/cpp/include/cuvs/neighbors/hnsw.hpp +++ b/cpp/include/cuvs/neighbors/hnsw.hpp @@ -16,6 +16,7 @@ #include #include #include +#include #include #include @@ -35,16 +36,37 @@ namespace graph_build_params = cuvs::neighbors::graph_build_params; * @brief Hierarchy for HNSW index when converting from CAGRA index * * NOTE: When the value is `NONE`, the HNSW index is built as a base-layer-only index. + * When the value is `CPU`, a full index is built with a CPU-constructed hierarchy. + * When the value is `GPU`, a full index is built with a GPU-constructed hierarchy. */ -enum class HnswHierarchy { - NONE, // base-layer-only index - CPU, // full index with CPU-built hierarchy - GPU // full index with GPU-built hierarchy +enum class HnswHierarchy { NONE, CPU, GPU }; + +/** + * @brief Output artifact format for an HNSW index + * + * `HNSWLIB` produces the standard hnswlib index format. `GRAPH_ONLY` stores the graph + * separately from vectors. The current implementation requires `HnswHierarchy::GPU` and + * disk-backed ACE. Load a graph-only artifact with the two-filename `deserialize` overload, which + * reads the vectors from a separate local dataset. + */ +enum class HnswOutputFormat { + HNSWLIB, + GRAPH_ONLY, }; +/** + * @code{.cpp} + * hnsw::index_params params; + * params.hierarchy = hnsw::HnswHierarchy::GPU; + * params.output_format = hnsw::HnswOutputFormat::GRAPH_ONLY; + * @endcode + */ struct index_params : cuvs::neighbors::index_params { /** Hierarchy build type for HNSW index when converting from CAGRA index */ HnswHierarchy hierarchy = HnswHierarchy::GPU; + /** Output artifact format. Graph-only output currently requires a GPU hierarchy and disk-backed + * ACE. */ + HnswOutputFormat output_format = HnswOutputFormat::HNSWLIB; /** Size of the candidate list during hierarchy construction when hierarchy is `CPU`*/ int ef_construction = 200; /** Number of host threads to use to construct hierarchy when hierarchy is `CPU` or `GPU`. @@ -81,9 +103,9 @@ struct index_params : cuvs::neighbors::index_params { * @endcode * * When ACE writes to disk, `build_dir` may already exist, but ACE's named CAGRA - * artifacts and `hnsw_index.bin` must not already exist. Simultaneous builds - * must use different directories. The HNSW output is published only after it - * is fully serialized; however, the complete build is not transactional: if + * artifacts and the selected HNSW output (`hnsw_index.bin` or `hnsw_index.cuvs`) must not already + * exist. Simultaneous builds must use different directories. The HNSW output is published only + * after it is fully serialized; however, the complete build is not transactional: if * HNSW conversion fails after CAGRA succeeds, the completed CAGRA artifacts * remain in the directory. */ @@ -145,9 +167,13 @@ struct index : cuvs::neighbors::index { * @param[in] dim dimensions of the training dataset * @param[in] metric distance metric to search. Supported metrics ("L2Expanded", "InnerProduct") * @param[in] hierarchy hierarchy used for upper HNSW layers + * @param[in] output_format output artifact format */ - index(int dim, cuvs::distance::DistanceType metric, HnswHierarchy hierarchy = HnswHierarchy::NONE) - : dim_{dim}, metric_{metric}, hierarchy_{hierarchy} + index(int dim, + cuvs::distance::DistanceType metric, + HnswHierarchy hierarchy = HnswHierarchy::NONE, + HnswOutputFormat output_format = HnswOutputFormat::HNSWLIB) + : dim_{dim}, metric_{metric}, hierarchy_{hierarchy}, output_format_{output_format} { } @@ -164,6 +190,8 @@ struct index : cuvs::neighbors::index { auto hierarchy() const -> HnswHierarchy { return hierarchy_; } + auto output_format() const -> HnswOutputFormat { return output_format_; } + /** @brief Set ef for search */ @@ -178,6 +206,7 @@ struct index : cuvs::neighbors::index { int dim_; cuvs::distance::DistanceType metric_; HnswHierarchy hierarchy_; + HnswOutputFormat output_format_; }; /** @@ -454,6 +483,10 @@ std::unique_ptr> build( * the format is not compatible with the original hnswlib. * 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The * serialized index is also compatible with the original hnswlib library. + * 3. `GPU`: The hierarchy is constructed on the GPU. + * When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. + * Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the + * local dataset. * * @param[in] res raft resources * @param[in] params hnsw index parameters @@ -490,6 +523,10 @@ std::unique_ptr> from_cagra( * the format is not compatible with the original hnswlib. * 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The * serialized index is also compatible with the original hnswlib library. + * 3. `GPU`: The hierarchy is constructed on the GPU. + * When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. + * Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the + * local dataset. * * @param[in] res raft resources * @param[in] params hnsw index parameters @@ -526,6 +563,10 @@ std::unique_ptr> from_cagra( * the format is not compatible with the original hnswlib. * 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The * serialized index is also compatible with the original hnswlib library. + * 3. `GPU`: The hierarchy is constructed on the GPU. + * When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. + * Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the + * local dataset. * * @param[in] res raft resources * @param[in] params hnsw index parameters @@ -562,6 +603,10 @@ std::unique_ptr> from_cagra( * the format is not compatible with the original hnswlib. * 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The * serialized index is also compatible with the original hnswlib library. + * 3. `GPU`: The hierarchy is constructed on the GPU. + * When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. + * Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the + * local dataset. * * @param[in] res raft resources * @param[in] params hnsw index parameters @@ -862,10 +907,12 @@ struct search_params : cuvs::neighbors::search_params { * @brief Search HNSW index constructed from a CAGRA index * NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is * `NONE`, as the format is not compatible with the original hnswlib. + * When `output_format` is `GRAPH_ONLY`, search uses the in-memory index reconstructed + * from the graph-only artifact produced by the two-filename `deserialize` overload. * * @param[in] res raft resources * @param[in] params configure the search - * @param[in] idx cagra index + * @param[in] idx HNSW index * @param[in] queries a host matrix view to a row-major matrix [n_queries, index->dim()] * @param[out] neighbors a host matrix view to the indices of the neighbors in the source dataset * [n_queries, k] @@ -906,10 +953,12 @@ void search(raft::resources const& res, * @brief Search HNSW index constructed from a CAGRA index * NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is * `NONE`, as the format is not compatible with the original hnswlib. + * When `output_format` is `GRAPH_ONLY`, search uses the in-memory index reconstructed + * from the graph-only artifact produced by the two-filename `deserialize` overload. * * @param[in] res raft resources * @param[in] params configure the search - * @param[in] idx cagra index + * @param[in] idx HNSW index * @param[in] queries a host matrix view to a row-major matrix [n_queries, index->dim()] * @param[out] neighbors a host matrix view to the indices of the neighbors in the source dataset * [n_queries, k] @@ -950,10 +999,12 @@ void search(raft::resources const& res, * @brief Search HNSWindex constructed from a CAGRA index * NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is * `NONE`, as the format is not compatible with the original hnswlib. + * When `output_format` is `GRAPH_ONLY`, search uses the in-memory index reconstructed + * from the graph-only artifact produced by the two-filename `deserialize` overload. * * @param[in] res raft resources * @param[in] params configure the search - * @param[in] idx cagra index + * @param[in] idx HNSW index * @param[in] queries a host matrix view to a row-major matrix [n_queries, index->dim()] * @param[out] neighbors a host matrix view to the indices of the neighbors in the source dataset * [n_queries, k] @@ -994,10 +1045,12 @@ void search(raft::resources const& res, * @brief Search HNSW index constructed from a CAGRA index * NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is * `NONE`, as the format is not compatible with the original hnswlib. + * When `output_format` is `GRAPH_ONLY`, search uses the in-memory index reconstructed + * from the graph-only artifact produced by the two-filename `deserialize` overload. * * @param[in] res raft resources * @param[in] params configure the search - * @param[in] idx cagra index + * @param[in] idx HNSW index * @param[in] queries a host matrix view to a row-major matrix [n_queries, index->dim()] * @param[out] neighbors a host matrix view to the indices of the neighbors in the source dataset * [n_queries, k] @@ -1039,7 +1092,7 @@ void search(raft::resources const& res, */ /** - * @defgroup hnsw_cpp_index_serialize Deserialize CAGRA index as hnswlib index + * @defgroup hnsw_cpp_index_serialize Serialize and deserialize HNSW indexes * @{ */ @@ -1049,10 +1102,12 @@ void search(raft::resources const& res, * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. + * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. + * Load it with the two-filename `deserialize` overload and a local dataset. * * @param[in] res raft resources - * @param[in] filename path to the file to save the serialized CAGRA index - * @param[in] idx cagra index + * @param[in] filename path to the serialized HNSW output + * @param[in] idx HNSW index * * Usage example: * @code{.cpp} @@ -1067,7 +1122,7 @@ void search(raft::resources const& res, * hnsw::index_params hnsw_params; * auto hnsw_index = hnsw::from_cagra(res, hnsw_params, index); * // Save the index - * hnsw::serialize(res, "index.bin", index); + * hnsw::serialize(res, "index.bin", *hnsw_index); * @endcode */ void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -1078,10 +1133,12 @@ void serialize(raft::resources const& res, const std::string& filename, const in * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. + * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. + * Load it with the two-filename `deserialize` overload and a local dataset. * * @param[in] res raft resources - * @param[in] filename path to the file to save the serialized CAGRA index - * @param[in] idx cagra index + * @param[in] filename path to the serialized HNSW output + * @param[in] idx HNSW index * * Usage example: * @code{.cpp} @@ -1096,7 +1153,7 @@ void serialize(raft::resources const& res, const std::string& filename, const in * hnsw::index_params hnsw_params; * auto hnsw_index = hnsw::from_cagra(res, hnsw_params, index); * // Save the index - * hnsw::serialize(res, "index.bin", index); + * hnsw::serialize(res, "index.bin", *hnsw_index); * @endcode */ void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -1107,10 +1164,12 @@ void serialize(raft::resources const& res, const std::string& filename, const in * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. + * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. + * Load it with the two-filename `deserialize` overload and a local dataset. * * @param[in] res raft resources - * @param[in] filename path to the file to save the serialized CAGRA index - * @param[in] idx cagra index + * @param[in] filename path to the serialized HNSW output + * @param[in] idx HNSW index * * Usage example: * @code{.cpp} @@ -1125,7 +1184,7 @@ void serialize(raft::resources const& res, const std::string& filename, const in * hnsw::index_params hnsw_params; * auto hnsw_index = hnsw::from_cagra(res, hnsw_params, index); * // Save the index - * hnsw::serialize(res, "index.bin", index); + * hnsw::serialize(res, "index.bin", *hnsw_index); * @endcode */ void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -1136,10 +1195,12 @@ void serialize(raft::resources const& res, const std::string& filename, const in * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. + * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. + * Load it with the two-filename `deserialize` overload and a local dataset. * * @param[in] res raft resources - * @param[in] filename path to the file to save the serialized CAGRA index - * @param[in] idx cagra index + * @param[in] filename path to the serialized HNSW output + * @param[in] idx HNSW index * * Usage example: * @code{.cpp} @@ -1154,21 +1215,22 @@ void serialize(raft::resources const& res, const std::string& filename, const in * hnsw::index_params hnsw_params; * auto hnsw_index = hnsw::from_cagra(res, hnsw_params, index); * // Save the index - * hnsw::serialize(res, "index.bin", index); + * hnsw::serialize(res, "index.bin", *hnsw_index); * @endcode */ void serialize(raft::resources const& res, const std::string& filename, const index& idx); /** - * @brief De-serialize a CAGRA index saved to a file as an hnswlib index + * @brief Deserialize an HNSWLIB index * NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. + * This overload loads `HNSWLIB` artifacts. Use the two-filename overload for graph-only artifacts. * * @param[in] res raft resources * @param[in] params hnsw index parameters - * @param[in] filename path to the file containing the serialized CAGRA index + * @param[in] filename path to the HNSWLIB artifact * @param[in] dim dimensions of the training dataset * @param[in] metric distance metric to search. Supported metrics ("L2Expanded", "InnerProduct") * @param[out] index hnsw index @@ -1186,13 +1248,14 @@ void serialize(raft::resources const& res, const std::string& filename, const in * hnsw::index_params hnsw_params; * auto hnsw_index = hnsw::from_cagra(res, hnsw_params, index); * // save HNSW index to a file - * hnsw::serialize(res, "index.bin", hnsw_index); - * // De-serialize the HNSW index - * index* hnsw_index = nullptr; - * hnsw::deserialize(res, hnsw_params, "index.bin", index->dim(), index->metric(), &hnsw_index); + * hnsw::serialize(res, "index.bin", *hnsw_index); + * // Deserialize the HNSW index + * hnsw::index* loaded_hnsw_index = nullptr; + * hnsw::deserialize( + * res, hnsw_params, "index.bin", index->dim(), index->metric(), &loaded_hnsw_index); * * // Delete index after use - * delete hnsw_index; + * delete loaded_hnsw_index; * @endcode */ void deserialize(raft::resources const& res, @@ -1203,15 +1266,16 @@ void deserialize(raft::resources const& res, index** index); /** - * @brief De-serialize a CAGRA index saved to a file as an hnswlib index + * @brief Deserialize an HNSWLIB index * NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. + * This overload loads `HNSWLIB` artifacts. Use the two-filename overload for graph-only artifacts. * * @param[in] res raft resources * @param[in] params hnsw index parameters - * @param[in] filename path to the file containing the serialized CAGRA index + * @param[in] filename path to the HNSWLIB artifact * @param[in] dim dimensions of the training dataset * @param[in] metric distance metric to search. Supported metrics ("L2Expanded", "InnerProduct") * @param[out] index hnsw index @@ -1229,13 +1293,14 @@ void deserialize(raft::resources const& res, * hnsw::index_params hnsw_params; * auto hnsw_index = hnsw::from_cagra(res, hnsw_params, index); * // save HNSW index to a file - * hnsw::serialize(res, "index.bin", hnsw_index); - * // De-serialize the HNSW index - * index* hnsw_index = nullptr; - * hnsw::deserialize(res, hnsw_params, "index.bin", index->dim(), index->metric(), &hnsw_index); + * hnsw::serialize(res, "index.bin", *hnsw_index); + * // Deserialize the HNSW index + * hnsw::index* loaded_hnsw_index = nullptr; + * hnsw::deserialize( + * res, hnsw_params, "index.bin", index->dim(), index->metric(), &loaded_hnsw_index); * * // Delete index after use - * delete hnsw_index; + * delete loaded_hnsw_index; * @endcode */ void deserialize(raft::resources const& res, @@ -1246,15 +1311,16 @@ void deserialize(raft::resources const& res, index** index); /** - * @brief De-serialize a CAGRA index saved to a file as an hnswlib index + * @brief Deserialize an HNSWLIB index * NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. + * This overload loads `HNSWLIB` artifacts. Use the two-filename overload for graph-only artifacts. * * @param[in] res raft resources * @param[in] params hnsw index parameters - * @param[in] filename path to the file containing the serialized CAGRA index + * @param[in] filename path to the HNSWLIB artifact * @param[in] dim dimensions of the training dataset * @param[in] metric distance metric to search. Supported metrics ("L2Expanded", "InnerProduct") * @param[out] index hnsw index @@ -1272,13 +1338,14 @@ void deserialize(raft::resources const& res, * hnsw::index_params hnsw_params; * auto hnsw_index = hnsw::from_cagra(res, hnsw_params, index); * // save HNSW index to a file - * hnsw::serialize(res, "index.bin", hnsw_index); - * // De-serialize the HNSW index - * index* hnsw_index = nullptr; - * hnsw::deserialize(res, hnsw_params, "index.bin", index->dim(), index->metric(), &hnsw_index); + * hnsw::serialize(res, "index.bin", *hnsw_index); + * // Deserialize the HNSW index + * hnsw::index* loaded_hnsw_index = nullptr; + * hnsw::deserialize( + * res, hnsw_params, "index.bin", index->dim(), index->metric(), &loaded_hnsw_index); * * // Delete index after use - * delete hnsw_index; + * delete loaded_hnsw_index; * @endcode */ void deserialize(raft::resources const& res, @@ -1289,15 +1356,16 @@ void deserialize(raft::resources const& res, index** index); /** - * @brief De-serialize a CAGRA index saved to a file as an hnswlib index + * @brief Deserialize an HNSWLIB index * NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. + * This overload loads `HNSWLIB` artifacts. Use the two-filename overload for graph-only artifacts. * * @param[in] res raft resources * @param[in] params hnsw index parameters - * @param[in] filename path to the file containing the serialized CAGRA index + * @param[in] filename path to the HNSWLIB artifact * @param[in] dim dimensions of the training dataset * @param[in] metric distance metric to search. Supported metrics ("L2Expanded", "InnerProduct") * @param[out] index hnsw index @@ -1315,13 +1383,14 @@ void deserialize(raft::resources const& res, * hnsw::index_params hnsw_params; * auto hnsw_index = hnsw::from_cagra(res, hnsw_params, index); * // save HNSW index to a file - * hnsw::serialize(res, "index.bin", hnsw_index); - * // De-serialize the HNSW index - * index* hnsw_index = nullptr; - * hnsw::deserialize(res, hnsw_params, "index.bin", index->dim(), index->metric(), &hnsw_index); + * hnsw::serialize(res, "index.bin", *hnsw_index); + * // Deserialize the HNSW index + * hnsw::index* loaded_hnsw_index = nullptr; + * hnsw::deserialize( + * res, hnsw_params, "index.bin", index->dim(), index->metric(), &loaded_hnsw_index); * * // Delete index after use - * delete hnsw_index; + * delete loaded_hnsw_index; * @endcode */ void deserialize(raft::resources const& res, @@ -1331,6 +1400,53 @@ void deserialize(raft::resources const& res, cuvs::distance::DistanceType metric, index** index); +/** + * @brief Deserialize a graph-only HNSW artifact and attach its dataset + * + * The graph artifact supplies the index dimensions, metric, and construction metadata. The + * attached dataset must have the recorded shape, but its element type may differ from the type used + * to construct the graph. The output pointer selects the attached dataset type. The loader accepts + * row-major `.npy` files and ANN benchmark binary files with a `[uint32 rows, uint32 cols]` header. + * Binary extensions must match the output index type: `.fbin` for `float`, `.f16bin` or + * `.fp16.fbin` for `half`, `.u8bin` for `uint8_t`, and `.i8bin` for `int8_t`. + * + * @param[in] res raft resources + * @param[in] graph_filename path to the graph-only HNSW artifact + * @param[in] dataset_filename path to the local dataset + * @param[out] index reconstructed HNSW index + * + * Usage example: + * @code{.cpp} + * using namespace cuvs::neighbors; + * hnsw::index* index = nullptr; + * hnsw::deserialize(res, "hnsw_index.cuvs", "dataset.fbin", &index); + * // ... search index ... + * delete index; + * @endcode + */ +void deserialize(raft::resources const& res, + const std::string& graph_filename, + const std::string& dataset_filename, + index** index); + +/** @overload */ +void deserialize(raft::resources const& res, + const std::string& graph_filename, + const std::string& dataset_filename, + index** index); + +/** @overload */ +void deserialize(raft::resources const& res, + const std::string& graph_filename, + const std::string& dataset_filename, + index** index); + +/** @overload */ +void deserialize(raft::resources const& res, + const std::string& graph_filename, + const std::string& dataset_filename, + index** index); + /** * @} */ diff --git a/cpp/include/cuvs/util/file_io.hpp b/cpp/include/cuvs/util/file_io.hpp index 84937945d1..e2e9d34026 100644 --- a/cpp/include/cuvs/util/file_io.hpp +++ b/cpp/include/cuvs/util/file_io.hpp @@ -8,6 +8,8 @@ #include #include +#include + #include #include #include @@ -330,14 +332,7 @@ std::pair create_numpy_file(const std::string& path, file_descriptor fd(path, flags, 0644); try { - // Build header - const auto dtype = raft::numpy_serializer::get_numpy_dtype(); - const bool fortran_order = false; - const raft::numpy_serializer::header_t header = {dtype, fortran_order, shape}; - - std::stringstream ss; - raft::numpy_serializer::write_header(ss, header); - std::string header_str = ss.str(); + std::string header_str = detail::make_numpy_header_string(shape); RAFT_EXPECTS((kNumpyDataAlignment & (kNumpyDataAlignment - 1)) == 0, "kNumpyDataAlignment must be a power of two"); diff --git a/cpp/include/cuvs/util/numpy_dtype.hpp b/cpp/include/cuvs/util/numpy_dtype.hpp new file mode 100644 index 0000000000..a06d9b27a8 --- /dev/null +++ b/cpp/include/cuvs/util/numpy_dtype.hpp @@ -0,0 +1,131 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include + +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +namespace CUVS_EXPORT cuvs { +namespace util { +namespace detail { + +/** + * @brief Return the NumPy dtype descriptor for a cuVS scalar type. + * + * @tparam T Scalar type to describe. + * + * @return A NumPy dtype descriptor string, for example `" +inline auto numpy_dtype_string() -> std::string +{ + if constexpr (std::is_same_v) { + return "().to_string(); + } +} + +/** + * @brief Build a complete NumPy v1 `.npy` header for a known dtype descriptor and shape. + * + * @param dtype NumPy dtype descriptor to serialize into the header dictionary, such as `"& shape) + -> std::string +{ + std::stringstream dict; + dict << "{'descr': '" << dtype << "', 'fortran_order': False, 'shape': ("; + for (size_t i = 0; i < shape.size(); ++i) { + if (i != 0) { dict << ", "; } + dict << shape[i]; + } + if (shape.size() == 1) { dict << ","; } + dict << "), }"; + + std::string header = dict.str(); + constexpr size_t preamble_size = + 6 + 2 + sizeof(uint16_t); // magic string, version, v1 header length + const size_t remainder = (preamble_size + header.size() + 1) % 16; + if (remainder != 0) { header.append(16 - remainder, ' '); } + header.push_back('\n'); + + RAFT_EXPECTS(header.size() <= std::numeric_limits::max(), + "NumPy v1 header is too large: %zu bytes", + header.size()); + + const auto header_len = static_cast(header.size()); + std::string result; + result.reserve(preamble_size + header.size()); + result.append("\x93NUMPY", 6); + result.push_back(1); + result.push_back(0); + result.push_back(static_cast(header_len & 0xff)); + result.push_back(static_cast((header_len >> 8) & 0xff)); + result.append(header); + return result; +} + +/** + * @brief Build a complete NumPy v1 `.npy` header for a cuVS scalar type and shape. + * + * @tparam T Scalar type of the serialized array. + * + * @param shape Row-major array extents to serialize in NumPy shape syntax. + * + * @return Serialized NumPy v1 header bytes, including magic string, version, little-endian header + * length, padded dictionary payload, and trailing newline. + * + * @throws raft::exception when the generated half-precision header is too large for NumPy v1. + * Other scalar types use RAFT's public header writer and propagate any exceptions from that path. + * + * @note `half` uses cuVS's explicit `" +inline auto make_numpy_header_string(const std::vector& shape) -> std::string +{ + if constexpr (std::is_same_v) { + return make_numpy_header_from_dtype(numpy_dtype_string(), shape); + } else { + const auto dtype = raft::numpy_serializer::get_numpy_dtype(); + const bool fortran_order = false; + const raft::numpy_serializer::header_t header = {dtype, fortran_order, shape}; + + std::stringstream ss; + raft::numpy_serializer::write_header(ss, header); + return ss.str(); + } +} + +} // namespace detail +} // namespace util +} // namespace CUVS_EXPORT cuvs diff --git a/cpp/src/neighbors/brute_force_serialize.cu b/cpp/src/neighbors/brute_force_serialize.cu index 53db313de2..c7aa4d1a37 100644 --- a/cpp/src/neighbors/brute_force_serialize.cu +++ b/cpp/src/neighbors/brute_force_serialize.cu @@ -28,7 +28,7 @@ void serialize(raft::resources const& handle, RAFT_LOG_DEBUG( "Saving brute force index, size %zu, dim %u", static_cast(index.size()), index.dim()); - auto dtype_string = raft::numpy_serializer::get_numpy_dtype().to_string(); + auto dtype_string = cuvs::util::detail::numpy_dtype_string(); dtype_string.resize(4); os << dtype_string; diff --git a/cpp/src/neighbors/detail/cagra/cagra_build.cuh b/cpp/src/neighbors/detail/cagra/cagra_build.cuh index d7a028f1cf..5fb60581e2 100644 --- a/cpp/src/neighbors/detail/cagra/cagra_build.cuh +++ b/cpp/src/neighbors/detail/cagra/cagra_build.cuh @@ -42,6 +42,7 @@ #include +#include #include #include #include @@ -65,6 +66,13 @@ namespace cuvs::neighbors::cagra::detail { constexpr double to_mib(size_t bytes) { return static_cast(bytes) / (1 << 20); } constexpr double to_gib(size_t bytes) { return static_cast(bytes) / (1 << 30); } +inline auto progress_step_10(size_t completed, size_t total) -> size_t +{ + if (total == 0) { return 100; } + const auto percent = (std::min(completed, total) * 100) / total; + return completed >= total ? 100 : (percent / 10) * 10; +} + class ace_disk_workspace { public: enum class artifact : size_t { @@ -206,7 +214,7 @@ void ace_get_partition_labels( const size_t min_samples = 100 * n_partitions; n_samples = std::max(n_samples, min_samples); n_samples = std::min(n_samples, dataset_size); - RAFT_LOG_DEBUG("ACE: n_samples: %lu", n_samples); + RAFT_LOG_DEBUG("ACE build: partition labeling uses %lu sampled vectors", n_samples); auto sample_db = raft::make_host_matrix(n_samples, dataset_dim); #pragma omp parallel for @@ -230,18 +238,10 @@ void ace_get_partition_labels( auto _sub_distances = raft::make_host_matrix(chunk_size, n_partitions); auto _sub_dataset_dev = raft::make_device_matrix(res, chunk_size, dataset_dim); auto _sub_distances_dev = raft::make_device_matrix(res, chunk_size, n_partitions); - size_t report_interval = dataset_size / 10; - report_interval = (report_interval / chunk_size) * chunk_size; - report_interval = std::max(report_interval, chunk_size); + size_t next_progress_percent = 10; for (size_t i_base = 0; i_base < dataset_size; i_base += chunk_size) { const size_t sub_dataset_size = std::min(chunk_size, dataset_size - i_base); - if (i_base % report_interval == 0) { - RAFT_LOG_INFO("ACE: Processing chunk %lu / %lu (%.1f%%)", - i_base, - dataset_size, - static_cast(100 * i_base) / dataset_size); - } auto sub_dataset = raft::make_host_matrix_view( _sub_dataset.data_handle(), sub_dataset_size, dataset_dim); @@ -298,6 +298,16 @@ void ace_get_partition_labels( #pragma omp atomic update partition_histogram(augmented_label, 1) += 1; } + + const auto completed_rows = i_base + sub_dataset_size; + const auto report_percent = progress_step_10(completed_rows, dataset_size); + if (report_percent >= next_progress_percent) { + RAFT_LOG_INFO("ACE build: partition labeling progress %zu%% (%zu/%zu rows)", + report_percent, + completed_rows, + dataset_size); + next_progress_percent = report_percent + 10; + } } } @@ -368,15 +378,15 @@ void ace_check_partition_sizes( if (total_count > 0 && total_count < very_small_threshold) { RAFT_LOG_WARN( - "ACE: Partition %lu is very small (%lu vectors, expected ~%.1f). This may affect graph " - "quality.", + "ACE build: partition %lu is very small (%lu vectors, expected ~%.1f); graph quality may " + "be affected", c, total_count, expected_avg_vectors); } else if (total_count > very_large_threshold) { RAFT_LOG_WARN( - "ACE: Partition %lu is very large (%lu vectors, expected ~%.1f, threshold: %lu). This may " - "indicate imbalance and can lead to memory issues in restricted environments.", + "ACE build: partition %lu is very large (%lu vectors, expected ~%.1f, threshold=%lu); " + "partition imbalance may increase memory pressure", c, total_count, expected_avg_vectors, @@ -641,11 +651,10 @@ void ace_reorder_and_store_dataset( static_cast(dataset.extent(1)), dataset_dim); - RAFT_LOG_DEBUG( - "ACE: Reordering and storing dataset to disk (%lu vectors, %lu dimensions, %lu partitions)", - dataset_size, - dataset_dim, - n_partitions); + RAFT_LOG_DEBUG("ACE build: reordering dataset to disk (rows=%lu dim=%lu partitions=%lu)", + dataset_size, + dataset_dim, + n_partitions); // Calculate total sizes for pre-allocation size_t total_core_vectors = 0; @@ -668,10 +677,10 @@ void ace_reorder_and_store_dataset( size_t reordered_file_size = total_core_vectors * vector_size; size_t augmented_file_size = total_augmented_vectors * vector_size; - RAFT_LOG_DEBUG("ACE: Reordered dataset: %lu core vectors (%.2f GiB)", + RAFT_LOG_DEBUG("ACE build: reordered dataset section rows=%lu size=%.2f GiB", total_core_vectors, reordered_file_size / (1024.0 * 1024.0 * 1024.0)); - RAFT_LOG_DEBUG("ACE: Augmented dataset: %lu secondary vectors (%.2f GiB)", + RAFT_LOG_DEBUG("ACE build: augmented dataset section rows=%lu size=%.2f GiB", total_augmented_vectors, augmented_file_size / (1024.0 * 1024.0 * 1024.0)); @@ -700,7 +709,7 @@ void ace_reorder_and_store_dataset( disk_write_size = std::min(disk_write_size, 64 * 1024 * 1024); size_t vectors_per_buffer = std::max(64, disk_write_size / vector_size); - RAFT_LOG_DEBUG("ACE: Reorder buffers: %lu vectors per buffer (%.2f MiB)", + RAFT_LOG_DEBUG("ACE build: reorder buffer rows=%lu size=%.2f MiB", vectors_per_buffer, to_mib(vectors_per_buffer * vector_size)); @@ -752,8 +761,8 @@ void ace_reorder_and_store_dataset( } }; - size_t vectors_processed = 0; - const size_t log_interval = std::max(dataset_size / 10, size_t(1)); + size_t vectors_processed = 0; + size_t next_progress_percent = 10; for (size_t i = 0; i < dataset_size; i++) { size_t core_partition = partition_labels(i, 0); size_t secondary_partition = partition_labels(i, 1); @@ -782,16 +791,18 @@ void ace_reorder_and_store_dataset( } vectors_processed++; - if (vectors_processed % log_interval == 0) { - RAFT_LOG_INFO("ACE: Processed %lu/%lu vectors (%.1f%%)", + const auto report_percent = progress_step_10(vectors_processed, dataset_size); + if (report_percent >= next_progress_percent) { + RAFT_LOG_INFO("ACE build: dataset reorder progress %zu%% (%zu/%zu rows)", + report_percent, vectors_processed, - dataset_size, - 100.0 * vectors_processed / dataset_size); + dataset_size); + next_progress_percent = report_percent + 10; } } // Flush all remaining buffers - RAFT_LOG_DEBUG("ACE: Flushing remaining buffers..."); + RAFT_LOG_DEBUG("ACE build: flushing remaining reorder buffers"); #pragma omp parallel sections { #pragma omp section @@ -821,12 +832,12 @@ void ace_reorder_and_store_dataset( elapsed_ms > 0 ? to_mib(total_bytes_written) / (elapsed_ms / 1000.0) : 0.0; RAFT_LOG_INFO( - "ACE: Dataset (%.2f GiB reordered, %.2f GiB augmented, %.2f GiB mapping) reordering completed " - "in %ld ms (%.1f MiB/s)", + "ACE build: dataset reorder completed in %ld ms (reordered=%.2f GiB augmented=%.2f GiB " + "mapping=%.2f GiB, %.1f MiB/s)", + elapsed_ms, reordered_file_size / (1024.0 * 1024.0 * 1024.0), augmented_file_size / (1024.0 * 1024.0 * 1024.0), mapping_file_size / (1024.0 * 1024.0 * 1024.0), - elapsed_ms, throughput_mb_s); } @@ -874,11 +885,11 @@ void ace_load_partition_dataset_from_disk( static_cast(sub_dataset.extent(1)), dataset_dim); - RAFT_LOG_DEBUG("ACE: Partition %lu: %lu core + %lu augmented = %lu total vectors", + RAFT_LOG_DEBUG("ACE build: partition %lu rows=%lu (core=%lu augmented=%lu)", partition_id, + total_size, core_size, - augmented_size, - core_size + augmented_size); + augmented_size); const size_t vector_size = dataset_dim * sizeof(T); const size_t core_file_offset = @@ -887,7 +898,7 @@ void ace_load_partition_dataset_from_disk( augmented_header_size + static_cast(augmented_partition_offsets(partition_id)) * vector_size; - RAFT_LOG_DEBUG("ACE: Core file offset: %lu bytes, Augmented file offset: %lu bytes", + RAFT_LOG_DEBUG("ACE build: core file offset=%lu bytes augmented file offset=%lu bytes", core_file_offset, augmented_file_offset); @@ -1053,7 +1064,7 @@ bool ace_check_use_disk_mode(raft::resources const& res, mem.total_size = mem.partition_labels_size + mem.id_mapping_size + mem.sub_dataset_size + mem.sub_graph_size + mem.cagra_graph_size + opt_host_ws_total; - RAFT_LOG_INFO("ACE: Estimated host memory required: %.2f GiB, available: %.2f GiB", + RAFT_LOG_INFO("ACE build: host memory estimate required=%.2f GiB available=%.2f GiB", to_gib(mem.total_size), to_gib(mem.available_host_memory)); @@ -1068,7 +1079,7 @@ bool ace_check_use_disk_mode(raft::resources const& res, auto configured_gpu_memory = static_cast(max_gpu_memory_gb.value() * (1ULL << 30)); if (actual_available_gpu_memory < configured_gpu_memory) { RAFT_LOG_WARN( - "ACE: Actual GPU memory (%zu GiB) is less than configured limit (%zu GiB). " + "ACE: Actual GPU memory (%.2f GiB) is less than configured limit (%.2f GiB). " "Using actual GPU memory.", to_gib(actual_available_gpu_memory), to_gib(configured_gpu_memory)); @@ -1099,24 +1110,19 @@ bool ace_check_use_disk_mode(raft::resources const& res, if (use_disk_mode) { RAFT_EXPECTS(!build_dir.empty(), "ACE build directory must not be empty"); } if (host_memory_limited && gpu_memory_limited) { - RAFT_LOG_INFO( - "ACE: Graph does not fit in host and GPU memory. Using disk-mode with temporary storage %s", - build_dir.c_str()); + RAFT_LOG_INFO("ACE build: graph does not fit in host or GPU memory; using disk mode at %s", + build_dir.c_str()); } else if (host_memory_limited) { - RAFT_LOG_INFO( - "ACE: Graph does not fit in host memory. Using disk-mode with temporary storage %s", - build_dir.c_str()); + RAFT_LOG_INFO("ACE build: graph does not fit in host memory; using disk mode at %s", + build_dir.c_str()); } else if (gpu_memory_limited) { - RAFT_LOG_INFO( - "ACE: Graph does not fit in GPU memory. Using disk-mode with temporary storage %s", - build_dir.c_str()); + RAFT_LOG_INFO("ACE build: graph does not fit in GPU memory; using disk mode at %s", + build_dir.c_str()); } else if (use_disk) { - RAFT_LOG_INFO( - "ACE: Graph fits in host and GPU memory but disk mode is forced. Using disk-mode with " - "temporary storage %s", - build_dir.c_str()); + RAFT_LOG_INFO("ACE build: graph fits in host and GPU memory but disk mode is forced; using %s", + build_dir.c_str()); } else { - RAFT_LOG_INFO("ACE: Graph fits in host and GPU memory. Using in-memory mode."); + RAFT_LOG_INFO("ACE build: graph fits in host and GPU memory; using in-memory mode"); } return use_disk_mode; @@ -1192,8 +1198,8 @@ void ace_validate_disk_mode_partitions(raft::resources const& res, disk_mode_host_required) { host_memory_insufficient = true; RAFT_LOG_WARN( - "ACE: Host memory insufficient for disk mode. Required: %.2f GiB, available: %.2f GiB. " - "Per-partition breakdown: dataset %.2f GiB, graph %.2f GiB, workspace %.2f GiB", + "ACE build: host memory is insufficient for disk mode: required=%.2f GiB available=%.2f " + "GiB (partition dataset=%.2f GiB graph=%.2f GiB workspace=%.2f GiB)", to_gib(disk_mode_host_required), to_gib(mem.available_host_memory), to_gib(mem.sub_dataset_size), @@ -1258,8 +1264,7 @@ void ace_validate_disk_mode_partitions(raft::resources const& res, size_t new_n_partitions = std::max(host_suggested_partitions, gpu_suggested_partitions); RAFT_LOG_WARN( - "ACE: Automatically increasing number of partitions from %zu to %zu to satisfy memory " - "constraints.%s%s", + "ACE build: increasing partitions from %zu to %zu to satisfy memory constraints.%s%s", original_n_partitions, new_n_partitions, host_memory_insufficient @@ -1288,8 +1293,8 @@ void ace_validate_disk_mode_partitions(raft::resources const& res, mem.sub_graph_size + mem.cagra_graph_size + new_opt_host_ws; RAFT_LOG_INFO( - "ACE: Updated per-partition memory estimates: dataset %.2f GiB, graph %.2f GiB, " - "host workspace %.2f GiB, GPU workspace %.2f GiB", + "ACE build: updated partition memory estimates dataset=%.2f GiB graph=%.2f GiB host " + "workspace=%.2f GiB GPU workspace=%.2f GiB", to_gib(mem.sub_dataset_size), to_gib(mem.sub_graph_size), to_gib(new_opt_host_ws), @@ -1346,15 +1351,15 @@ auto build_ace(raft::resources const& res, const index_params& params, DatasetVi size_t dataset_size = dataset.n_rows(); size_t dataset_dim = dataset.dim(); - RAFT_EXPECTS(dataset_size > 0, "ACE: Dataset must not be empty"); + RAFT_EXPECTS(dataset_size > 0, "ACE build: dataset must not be empty"); if (dataset_size < 1000) { - RAFT_LOG_WARN("ACE: Very small dataset size (%zu), consider using regular CAGRA build instead.", + RAFT_LOG_WARN("ACE build: very small dataset size (%zu); regular CAGRA may be simpler", dataset_size); } - RAFT_EXPECTS(dataset_dim > 0, "ACE: Dataset dimension must be greater than 0"); + RAFT_EXPECTS(dataset_dim > 0, "ACE build: dataset dimension must be greater than 0"); RAFT_EXPECTS(params.intermediate_graph_degree > 0, - "ACE: Intermediate graph degree must be greater than 0"); - RAFT_EXPECTS(params.graph_degree > 0, "ACE: Graph degree must be greater than 0"); + "ACE build: intermediate graph degree must be greater than 0"); + RAFT_EXPECTS(params.graph_degree > 0, "ACE build: graph degree must be greater than 0"); size_t n_partitions = ace_resolve_partition_count(npartitions); @@ -1365,17 +1370,22 @@ auto build_ace(raft::resources const& res, const index_params& params, DatasetVi n_partitions = dataset_size / min_required_per_partition; if (n_partitions < 2) { RAFT_LOG_WARN( - "ACE: Reduced number of partitions to the minimum of 2 to avoid tiny partitions. Consider " - "using regular CAGRA build instead."); + "ACE build: reduced partitions to the minimum of 2 to avoid tiny partitions; regular " + "CAGRA may be simpler"); n_partitions = 2; } else { - RAFT_LOG_WARN("ACE: Reduced number of partitions to %zu to avoid tiny partitions", - n_partitions); + RAFT_LOG_WARN("ACE build: reduced partitions to %zu to avoid tiny partitions", n_partitions); } } auto total_start = std::chrono::high_resolution_clock::now(); - RAFT_LOG_INFO("ACE: Starting partitioned CAGRA build with %zu partitions", n_partitions); + RAFT_LOG_INFO( + "ACE build: start rows=%zu dim=%zu partitions=%zu graph_degree=%zu intermediate_degree=%zu", + dataset_size, + dataset_dim, + n_partitions, + static_cast(params.graph_degree), + static_cast(params.intermediate_graph_degree)); size_t intermediate_degree = params.intermediate_graph_degree; size_t graph_degree = params.graph_degree; @@ -1452,7 +1462,7 @@ auto build_ace(raft::resources const& res, const index_params& params, DatasetVi workspace.mark_artifact_created(ace_disk_workspace::artifact::cagra_graph); RAFT_LOG_DEBUG( - "ACE: Wrote numpy headers (reordered: %zu, augmented: %zu, mapping: %zu, graph: %zu bytes)", + "ACE build: wrote numpy headers (reordered=%zu augmented=%zu mapping=%zu graph=%zu bytes)", reordered_header_size, augmented_header_size, mapping_header_size, @@ -1487,11 +1497,9 @@ auto build_ace(raft::resources const& res, const index_params& params, DatasetVi auto partition_elapsed = std::chrono::duration_cast(partition_end - partition_start) .count(); - RAFT_LOG_INFO( - "ACE: Partition labeling completed in %ld ms (min_partition_size: " - "%lu)", - partition_elapsed, - min_partition_size); + RAFT_LOG_INFO("ACE build: partition labeling completed in %ld ms (min_partition_size=%lu)", + partition_elapsed, + min_partition_size); // Create vector lists for each partition auto vectorlist_start = std::chrono::high_resolution_clock::now(); @@ -1516,7 +1524,7 @@ auto build_ace(raft::resources const& res, const index_params& params, DatasetVi auto vectorlist_elapsed = std::chrono::duration_cast(vectorlist_end - vectorlist_start) .count(); - RAFT_LOG_INFO("ACE: Vector list creation completed in %ld ms", vectorlist_elapsed); + RAFT_LOG_INFO("ACE build: partition mapping completed in %ld ms", vectorlist_elapsed); // Reorder the dataset based on partitions and store to disk. Uses write buffers to improve // performance. @@ -1548,7 +1556,7 @@ auto build_ace(raft::resources const& res, const index_params& params, DatasetVi // Process each partition auto partition_processing_start = std::chrono::high_resolution_clock::now(); for (size_t partition_id = 0; partition_id < n_partitions; partition_id++) { - RAFT_LOG_DEBUG("ACE: Processing partition %lu/%lu", partition_id + 1, n_partitions); + RAFT_LOG_DEBUG("ACE build: processing partition %lu/%lu", partition_id + 1, n_partitions); auto start = std::chrono::high_resolution_clock::now(); // Extract vectors for this partition @@ -1557,10 +1565,11 @@ auto build_ace(raft::resources const& res, const index_params& params, DatasetVi size_t sub_dataset_size = core_sub_dataset_size + augmented_sub_dataset_size; if (sub_dataset_size == 0) { - RAFT_LOG_WARN("ACE: Skipping empty partition %lu", partition_id); + RAFT_LOG_WARN("ACE build: Skipping empty partition %lu", partition_id); continue; } - RAFT_LOG_DEBUG("ACE: Sub-dataset size: %lu (%lu + %lu)", + RAFT_LOG_DEBUG("ACE build: partition %lu rows=%lu (core=%lu augmented=%lu)", + partition_id, sub_dataset_size, core_sub_dataset_size, augmented_sub_dataset_size); @@ -1740,12 +1749,12 @@ auto build_ace(raft::resources const& res, const index_params& params, DatasetVi double write_throughput = write_elapsed > 0 ? to_mib(graph_bytes) / (write_elapsed / 1000.0) : 0.0; RAFT_LOG_INFO( - "ACE: Partition %4lu (%8lu + %8lu) completed in %6ld ms: read %6ld ms (%7.1f MiB/s), " - "optimize %6ld ms, adjust %6ld ms, write %6ld ms (%7.1f MiB/s)", + "ACE build: partition %4lu completed in %6ld ms (core=%8lu augmented=%8lu read=%6ld ms " + "%.1f MiB/s optimize=%6ld ms adjust=%6ld ms write=%6ld ms %.1f MiB/s)", partition_id, + elapsed_ms, core_sub_dataset_size, augmented_sub_dataset_size, - elapsed_ms, read_elapsed, read_throughput, optimize_elapsed, @@ -1758,7 +1767,7 @@ auto build_ace(raft::resources const& res, const index_params& params, DatasetVi auto partition_processing_elapsed = std::chrono::duration_cast( partition_processing_end - partition_processing_start) .count(); - RAFT_LOG_INFO("ACE: All partition processing completed in %ld ms (%zu partitions)", + RAFT_LOG_INFO("ACE build: partition graph build completed in %ld ms (partitions=%zu)", partition_processing_elapsed, n_partitions); @@ -1768,7 +1777,7 @@ auto build_ace(raft::resources const& res, const index_params& params, DatasetVi const std::string augmented_dataset_path = build_dir + "/augmented_dataset.npy"; if (std::filesystem::exists(augmented_dataset_path)) { std::filesystem::remove(augmented_dataset_path); - RAFT_LOG_INFO("ACE: Removed augmented dataset file to save disk space"); + RAFT_LOG_INFO("ACE build: removed temporary augmented dataset"); } } @@ -1786,30 +1795,29 @@ auto build_ace(raft::resources const& res, const index_params& params, DatasetVi idx.update_graph(res, std::move(graph_fd)); idx.update_mapping(res, std::move(mapping_fd)); - RAFT_LOG_INFO( - "ACE: Set disk storage at %s (dataset shape [%zu, %zu], graph shape [%zu, %zu])", - build_dir.c_str(), - idx.size(), - idx.dim(), - idx.size(), - idx.graph_degree()); + RAFT_LOG_INFO("ACE build: disk artifacts ready at %s (dataset=[%zu,%zu] graph=[%zu,%zu])", + build_dir.c_str(), + idx.size(), + idx.dim(), + idx.size(), + idx.graph_degree()); } auto index_creation_end = std::chrono::high_resolution_clock::now(); auto index_creation_elapsed = std::chrono::duration_cast( index_creation_end - index_creation_start) .count(); - RAFT_LOG_INFO("ACE: Final index creation completed in %ld ms", index_creation_elapsed); + RAFT_LOG_INFO("ACE build: final index initialized in %ld ms", index_creation_elapsed); auto total_end = std::chrono::high_resolution_clock::now(); auto total_elapsed = std::chrono::duration_cast(total_end - total_start).count(); - RAFT_LOG_INFO("ACE: Partitioned CAGRA build completed in %ld ms total", total_elapsed); + RAFT_LOG_INFO("ACE build: completed in %ld ms", total_elapsed); workspace.commit(); return std::move(idx); } catch (const std::exception& e) { - RAFT_LOG_ERROR("ACE: Build failed with exception: %s", e.what()); + RAFT_LOG_ERROR("ACE build: failed with exception: %s", e.what()); workspace.cleanup(); throw; } catch (...) { diff --git a/cpp/src/neighbors/detail/cagra/cagra_serialize.cuh b/cpp/src/neighbors/detail/cagra/cagra_serialize.cuh index 6a59e80756..438303a607 100644 --- a/cpp/src/neighbors/detail/cagra/cagra_serialize.cuh +++ b/cpp/src/neighbors/detail/cagra/cagra_serialize.cuh @@ -119,7 +119,7 @@ void serialize(raft::resources const& res, if (include_dataset) { dataset_kind = serialized_dataset_kind_for_view(); } } - std::string dtype_string = raft::numpy_serializer::get_numpy_dtype().to_string(); + std::string dtype_string = cuvs::util::detail::numpy_dtype_string(); dtype_string.resize(4); os << dtype_string; diff --git a/cpp/src/neighbors/detail/hnsw.hpp b/cpp/src/neighbors/detail/hnsw.hpp index d39d525f80..19c0afd820 100644 --- a/cpp/src/neighbors/detail/hnsw.hpp +++ b/cpp/src/neighbors/detail/hnsw.hpp @@ -7,6 +7,7 @@ #include "../../core/nvtx.hpp" #include "../../core/omp_wrapper.hpp" +#include "hnsw_layered_format.hpp" #include #include @@ -28,27 +29,40 @@ #include +#include +#include #include +#include #include +#include +#include #include #include #include #include #include +#include #include #include #include #include +#include #include +#include +#include #include #include +#include #include +#include +#include namespace cuvs::neighbors::hnsw::detail { -class exclusive_hnsw_output_file { +// Writes to a sibling temp path and publishes with link(2) only after the write succeeds. +class exclusive_hnsw_temp_file { public: - explicit exclusive_hnsw_output_file(std::filesystem::path output_path) + explicit exclusive_hnsw_temp_file(std::filesystem::path output_path) : output_path_{std::move(output_path)} { std::string temporary_path = output_path_.string() + ".tmp.XXXXXX"; @@ -68,32 +82,18 @@ class exclusive_hnsw_output_file { error, strerror(error)); } - - try { - stream_ = std::make_unique(temporary_path_); - } catch (...) { - cleanup(); - throw; - } } - exclusive_hnsw_output_file(const exclusive_hnsw_output_file&) = delete; - exclusive_hnsw_output_file& operator=(const exclusive_hnsw_output_file&) = delete; + exclusive_hnsw_temp_file(const exclusive_hnsw_temp_file&) = delete; + exclusive_hnsw_temp_file& operator=(const exclusive_hnsw_temp_file&) = delete; - ~exclusive_hnsw_output_file() - { - stream_.reset(); - cleanup(); - } + ~exclusive_hnsw_temp_file() { cleanup(); } - std::ostream& stream() { return *stream_; } + [[nodiscard]] const std::filesystem::path& output_path() const { return output_path_; } + [[nodiscard]] const std::string& temporary_path() const { return temporary_path_; } void publish() { - stream_->close(); - RAFT_EXPECTS(*stream_, "Error writing output %s", output_path_.c_str()); - stream_.reset(); - if (::link(temporary_path_.c_str(), output_path_.c_str()) != 0) { const int error = errno; RAFT_FAIL("Cannot publish HNSW index %s (errno: %d, %s)", @@ -111,6 +111,33 @@ class exclusive_hnsw_output_file { std::filesystem::path output_path_; std::string temporary_path_; +}; + +class exclusive_hnsw_output_file { + public: + explicit exclusive_hnsw_output_file(std::filesystem::path output_path) + : temp_{std::move(output_path)} + { + stream_ = std::make_unique(temp_.temporary_path()); + } + + exclusive_hnsw_output_file(const exclusive_hnsw_output_file&) = delete; + exclusive_hnsw_output_file& operator=(const exclusive_hnsw_output_file&) = delete; + + ~exclusive_hnsw_output_file() { stream_.reset(); } + + std::ostream& stream() { return *stream_; } + + void publish() + { + stream_->close(); + RAFT_EXPECTS(*stream_, "Error writing output %s", temp_.output_path().c_str()); + stream_.reset(); + temp_.publish(); + } + + private: + exclusive_hnsw_temp_file temp_; std::unique_ptr stream_; }; @@ -196,9 +223,13 @@ struct index_impl : index { * @param[in] dim dimensions of the training dataset * @param[in] metric distance metric to search. Supported metrics ("L2Expanded", "InnerProduct") * @param[in] hierarchy hierarchy used for upper HNSW layers + * @param[in] output_format output artifact format */ - index_impl(int dim, cuvs::distance::DistanceType metric, HnswHierarchy hierarchy) - : index{dim, metric, hierarchy} + index_impl(int dim, + cuvs::distance::DistanceType metric, + HnswHierarchy hierarchy, + HnswOutputFormat output_format = HnswOutputFormat::HNSWLIB) + : index{dim, metric, hierarchy, output_format} { if (metric == cuvs::distance::DistanceType::InnerProduct) { space_ = std::make_unique::type>>(dim); @@ -287,6 +318,9 @@ struct index_impl : index { RAFT_LOG_INFO("Loading HNSW index from disk: %s", filepath.c_str()); try { + RAFT_EXPECTS(this->output_format() != HnswOutputFormat::GRAPH_ONLY, + "Layered HNSW indexes must be loaded with the two-filename hnsw::deserialize " + "overload so a local dataset can be provided."); appr_alg_ = std::make_unique::type>>( space_.get(), filepath); if (this->hierarchy() == HnswHierarchy::NONE) { appr_alg_->base_layer_only = true; } @@ -465,6 +499,830 @@ void all_neighbors_graph(raft::resources const& res, } } +/** Host-side plan that groups dataset rows by their maximum HNSW level. + * + * `row_ids_by_level[position]` maps a level-grouped position to the original dataset row ID. Rows + * whose maximum level is 0 come first, followed by rows whose maximum level is 1, and so on. Within + * each level, rows retain their original order. `position_by_row_id[row_id]` is the inverse + * mapping. `levels[row_id]` stores the maximum level assigned to that original row. The large + * row-sized arrays use RAFT host vectors so resource dry-run includes them in host-memory + * estimates. + */ +struct hnsw_level_plan { + size_t n_rows = 0; + std::vector hist; + // Bucket-end offsets after row ordering. Level L starts at offsets[L - 1]. + std::vector offsets; + raft::host_vector row_ids_by_level; + raft::host_vector position_by_row_id; + raft::host_vector levels; + + explicit hnsw_level_plan(raft::resources const& res) + : row_ids_by_level{raft::make_host_vector(res, 0)}, + position_by_row_id{raft::make_host_vector(res, 0)}, + levels{raft::make_host_vector(res, 0)} + { + } + + hnsw_level_plan(raft::resources const& res, + size_t n_rows, + raft::host_vector&& levels, + bool build_position_by_row_id) + : n_rows{n_rows}, + row_ids_by_level{raft::make_host_vector(res, n_rows)}, + position_by_row_id{raft::make_host_vector( + res, build_position_by_row_id ? n_rows : size_t{0})}, + levels{std::move(levels)} + { + } + + [[nodiscard]] auto max_level() const -> int + { + return hist.empty() ? 0 : static_cast(hist.size() - 1); + } + + [[nodiscard]] auto promoted_count() const -> size_t + { + return hist.empty() ? 0 : n_rows - hist[0]; + } +}; + +inline auto hnsw_upper_layer_degree(size_t row_count, size_t M) -> size_t +{ + return row_count > M ? M : row_count > 0 ? row_count - 1 : 0; +} + +inline auto elapsed_ms_since(std::chrono::steady_clock::time_point start) -> int64_t +{ + return std::chrono::duration_cast(std::chrono::steady_clock::now() - + start) + .count(); +} + +inline auto elapsed_ms(std::chrono::steady_clock::duration elapsed) -> double +{ + return std::chrono::duration(elapsed).count(); +} + +inline auto to_gib(size_t bytes) -> double +{ + return static_cast(bytes) / (1024.0 * 1024.0 * 1024.0); +} + +inline auto throughput_gib_per_s(size_t bytes, int64_t elapsed_ms) -> double +{ + return elapsed_ms > 0 ? to_gib(bytes) / (elapsed_ms / 1000.0) : 0.0; +} + +inline auto throughput_gib_per_s(size_t bytes, std::chrono::steady_clock::duration elapsed) + -> double +{ + const auto elapsed_s = std::chrono::duration(elapsed).count(); + return elapsed_s > 0.0 ? to_gib(bytes) / elapsed_s : 0.0; +} + +inline auto progress_step_10(size_t completed, size_t total) -> size_t +{ + if (total == 0) { return 100; } + const auto percent = (std::min(completed, total) * 100) / total; + return completed >= total ? 100 : (percent / 10) * 10; +} + +inline auto make_hnsw_level_plan_from_levels(raft::resources const& res, + size_t n_rows, + raft::host_vector&& levels, + bool build_position_by_row_id, + const char* log_prefix = nullptr) -> hnsw_level_plan +{ + RAFT_EXPECTS(n_rows > 0, "HNSW hierarchy requires at least one row"); + RAFT_EXPECTS(levels.size() == n_rows, + "HNSW level count (%zu) must match row count (%zu)", + levels.size(), + n_rows); + hnsw_level_plan plan{res, n_rows, std::move(levels), build_position_by_row_id}; + + for (size_t i = 0; i < n_rows; ++i) { + const auto level = plan.levels(i); + while (static_cast(level) >= plan.hist.size()) { + plan.hist.push_back(0); + } + plan.hist[level]++; + } + + plan.offsets.resize(plan.hist.size() + 1, 0); + for (size_t i = 0; i < plan.hist.size() - 1; ++i) { + plan.offsets[i + 1] = plan.offsets[i] + plan.hist[i]; + if (log_prefix != nullptr) { + RAFT_LOG_INFO("%s%zu : %zu", log_prefix, i + 1, n_rows - plan.offsets[i + 1]); + } + } + + for (size_t i = 0; i < n_rows; ++i) { + const auto level = static_cast(plan.levels(i)); + if (build_position_by_row_id) { plan.position_by_row_id(i) = plan.offsets[level]; } + plan.row_ids_by_level(plan.offsets[level]++) = i; + } + + return plan; +} + +template +auto make_random_hnsw_level_plan(raft::resources const& res, + size_t n_rows, + HnswAlgo& appr_algo, + const char* log_prefix) -> hnsw_level_plan +{ + auto levels = raft::make_host_vector(res, n_rows); + for (int64_t i = 0; i < static_cast(n_rows); ++i) { + const auto pt_level = appr_algo.getRandomLevel(appr_algo.mult_); + RAFT_EXPECTS(pt_level <= std::numeric_limits::max(), + "HNSW serialization only supports levels up to %u", + static_cast(std::numeric_limits::max())); + levels(i) = static_cast(pt_level); + } + + return make_hnsw_level_plan_from_levels(res, n_rows, std::move(levels), true, log_prefix); +} + +template +void build_hnsw_upper_layer_graphs( + raft::resources const& res, + raft::host_matrix_view promoted_dataset, + const hnsw_level_plan& plan, + size_t M, + cuvs::distance::DistanceType metric, + Callback&& callback) +{ + static_assert(std::is_same_v, + "HNSW upper-layer graph construction expects uint32_t neighbor ids"); + + const auto dim = static_cast(promoted_dataset.extent(1)); + for (size_t pt_level = 1; pt_level < plan.hist.size(); ++pt_level) { + const auto start_idx = plan.offsets[pt_level - 1]; + const auto removed_rows = start_idx - plan.offsets[0]; + const auto row_count = plan.n_rows - start_idx; + const auto neighbor_size = hnsw_upper_layer_degree(row_count, M); + + RAFT_LOG_INFO("Layered HNSW artifact: computing upper-layer neighbors level=%zu", pt_level); + auto host_neighbors = raft::make_host_matrix( + static_cast(row_count), static_cast(neighbor_size)); + + if (neighbor_size > 0) { + auto layer_dataset_view = raft::make_host_matrix_view( + promoted_dataset.data_handle() + removed_rows * dim, + static_cast(row_count), + promoted_dataset.extent(1)); + all_neighbors_graph(res, layer_dataset_view, host_neighbors.view(), metric); + } + + callback(pt_level, start_idx, host_neighbors); + } +} + +struct layered_hnsw_layer_info { + size_t level = 0; + size_t row_count = 0; + size_t degree = 0; + size_t node_offset = 0; + size_t link_offset = 0; +}; + +struct layered_hnsw_file_metadata { + size_t n_rows = 0; + size_t dim = 0; + size_t M = 0; + size_t maxM = 0; + size_t maxM0 = 0; + size_t ef_construction = 0; + size_t base_degree = 0; + size_t levels_bytes = 0; + size_t base_nodes_bytes = 0; + size_t base_link_row_bytes = 0; + size_t base_links_bytes = 0; + size_t upper_nodes_count = 0; + size_t upper_nodes_bytes = 0; + size_t upper_link_row_bytes = 0; + size_t upper_links_bytes = 0; + double mult = 0.0; + int maxlevel = 0; + int enterpoint_node = 0; + std::vector layers; +}; + +inline auto align_up(size_t value, size_t alignment) -> size_t +{ + return ((value + alignment - 1) / alignment) * alignment; +} + +template +constexpr auto layered_dtype_code() -> layered_hnsw_dtype +{ + if constexpr (std::is_same_v) { + return layered_hnsw_dtype::float32; + } else if constexpr (std::is_same_v) { + return layered_hnsw_dtype::float16; + } else if constexpr (std::is_same_v) { + return layered_hnsw_dtype::uint8; + } else if constexpr (std::is_same_v) { + return layered_hnsw_dtype::int8; + } else { + return layered_hnsw_dtype::unknown; + } +} + +inline auto is_supported_layered_dtype(layered_hnsw_dtype dtype) -> bool +{ + switch (dtype) { + case layered_hnsw_dtype::float32: + case layered_hnsw_dtype::float16: + case layered_hnsw_dtype::uint8: + case layered_hnsw_dtype::int8: return true; + default: return false; + } +} + +inline auto metric_name(cuvs::distance::DistanceType metric) -> const char* +{ + switch (metric) { + case cuvs::distance::DistanceType::L2Expanded: return "L2Expanded"; + case cuvs::distance::DistanceType::InnerProduct: return "InnerProduct"; + default: return "Unknown"; + } +} + +template +auto make_layered_hnsw_header(const layered_hnsw_file_metadata& metadata, + cuvs::distance::DistanceType metric) -> layered_hnsw_file_header +{ + layered_hnsw_file_header header{}; + std::strncpy(header.magic, layered_hnsw_magic, sizeof(header.magic) - 1); + header.version = layered_hnsw_version; + header.construction_dtype = static_cast(layered_dtype_code()); + header.metric = static_cast(metric); + header.num_layers = static_cast(metadata.layers.size()); + header.n_rows = metadata.n_rows; + header.dim = metadata.dim; + header.M = metadata.M; + header.maxM = metadata.maxM; + header.maxM0 = metadata.maxM0; + header.ef_construction = metadata.ef_construction; + header.base_degree = metadata.base_degree; + header.levels_bytes = metadata.levels_bytes; + header.base_nodes_bytes = metadata.base_nodes_bytes; + header.base_link_row_bytes = metadata.base_link_row_bytes; + header.base_links_bytes = metadata.base_links_bytes; + header.upper_nodes_count = metadata.upper_nodes_count; + header.upper_nodes_bytes = metadata.upper_nodes_bytes; + header.upper_link_row_bytes = metadata.upper_link_row_bytes; + header.upper_links_bytes = metadata.upper_links_bytes; + header.mult = metadata.mult; + header.maxlevel = metadata.maxlevel; + header.enterpoint_node = metadata.enterpoint_node; + return header; +} + +inline auto layered_hnsw_metadata_from_header(const layered_hnsw_file_header& header) + -> layered_hnsw_file_metadata +{ + layered_hnsw_file_metadata metadata; + metadata.n_rows = header.n_rows; + metadata.dim = header.dim; + metadata.M = header.M; + metadata.maxM = header.maxM; + metadata.maxM0 = header.maxM0; + metadata.ef_construction = header.ef_construction; + metadata.base_degree = header.base_degree; + metadata.levels_bytes = header.levels_bytes; + metadata.base_nodes_bytes = header.base_nodes_bytes; + metadata.base_link_row_bytes = header.base_link_row_bytes; + metadata.base_links_bytes = header.base_links_bytes; + metadata.upper_nodes_count = header.upper_nodes_count; + metadata.upper_nodes_bytes = header.upper_nodes_bytes; + metadata.upper_link_row_bytes = header.upper_link_row_bytes; + metadata.upper_links_bytes = header.upper_links_bytes; + metadata.mult = header.mult; + metadata.maxlevel = header.maxlevel; + metadata.enterpoint_node = header.enterpoint_node; + return metadata; +} + +struct npy_file { + cuvs::util::file_descriptor fd; + size_t header_size = 0; + std::vector shape; + std::string dtype; + bool fortran_order = false; +}; + +inline auto open_npy_file(const std::string& path) -> npy_file +{ + std::ifstream stream(path, std::ios::binary); + RAFT_EXPECTS(stream.good(), "Failed to open numpy file: %s", path.c_str()); + auto header = raft::numpy_serializer::read_header(stream); + auto header_size = static_cast(stream.tellg()); + auto fd = cuvs::util::file_descriptor(path, O_RDONLY); + return {std::move(fd), header_size, header.shape, header.dtype.to_string(), header.fortran_order}; +} + +template +inline void validate_npy_file(const npy_file& file, const std::string& path, const char* name) +{ + const auto expected_dtype = cuvs::util::detail::numpy_dtype_string(); + RAFT_EXPECTS(file.dtype == expected_dtype, + "%s dtype (%s) does not match expected dtype (%s): %s", + name, + file.dtype.c_str(), + expected_dtype.c_str(), + path.c_str()); + RAFT_EXPECTS( + !file.fortran_order, "%s must be row-major, got Fortran-order layout: %s", name, path.c_str()); +} + +inline auto ends_with(const std::string& value, const std::string& suffix) -> bool +{ + return value.size() >= suffix.size() && + value.compare(value.size() - suffix.size(), suffix.size(), suffix) == 0; +} + +inline void copy_file_overwrite(const std::filesystem::path& src, const std::filesystem::path& dst) +{ + std::error_code ec; + if (src == dst || + (std::filesystem::exists(dst, ec) && std::filesystem::equivalent(src, dst, ec))) { + return; + } + if (!dst.parent_path().empty()) { std::filesystem::create_directories(dst.parent_path()); } + std::filesystem::copy_file(src, dst, std::filesystem::copy_options::overwrite_existing); +} + +template +inline auto open_layered_dataset_file(const std::string& path) -> npy_file +{ + if (std::filesystem::path(path).extension() == ".npy") { return open_npy_file(path); } + + std::ifstream stream(path, std::ios::binary); + RAFT_EXPECTS(stream.good(), "Failed to open dataset file: %s", path.c_str()); + std::array header{}; + stream.read(reinterpret_cast(header.data()), sizeof(uint32_t) * header.size()); + RAFT_EXPECTS(stream.gcount() == static_cast(sizeof(uint32_t) * header.size()), + "Failed to read ANN benchmark binary dataset header: %s", + path.c_str()); + + const auto ext = std::filesystem::path(path).extension().string(); + if constexpr (std::is_same_v) { + RAFT_EXPECTS(ext == ".fbin" && !ends_with(path, ".fp16.fbin"), + "Expected a .fbin dataset for float layered HNSW load: %s", + path.c_str()); + } else if constexpr (std::is_same_v) { + RAFT_EXPECTS(ext == ".f16bin" || ends_with(path, ".fp16.fbin"), + "Expected a .f16bin or .fp16.fbin dataset for half layered HNSW load: %s", + path.c_str()); + } else if constexpr (std::is_same_v) { + RAFT_EXPECTS( + ext == ".u8bin", "Expected a .u8bin dataset for uint8 layered HNSW load: %s", path.c_str()); + } else if constexpr (std::is_same_v) { + RAFT_EXPECTS( + ext == ".i8bin", "Expected a .i8bin dataset for int8 layered HNSW load: %s", path.c_str()); + } + + auto fd = cuvs::util::file_descriptor(path, O_RDONLY); + return {std::move(fd), + sizeof(uint32_t) * header.size(), + {header[0], header[1]}, + cuvs::util::detail::numpy_dtype_string(), + false}; +} + +template +void write_layered_base_links_from_disk(const CagraIndexT& index_, + const cuvs::util::file_descriptor& output_fd, + size_t base_nodes_offset, + size_t base_links_offset, + size_t base_link_row_bytes, + size_t maxM0) +{ + using IdxT = typename CagraIndexT::index_type; + static_assert(std::is_same_v, + "Layered HNSW artifacts store topology ids as uint32_t"); + const auto& graph_fd_opt = index_.graph_fd(); + const auto& mapping_fd_opt = index_.mapping_fd(); + RAFT_EXPECTS(graph_fd_opt.has_value() && graph_fd_opt->is_valid(), + "Graph file descriptor is not available"); + RAFT_EXPECTS(mapping_fd_opt.has_value() && mapping_fd_opt->is_valid(), + "Mapping file descriptor is not available"); + + const auto graph_path = graph_fd_opt->get_path(); + const auto mapping_path = mapping_fd_opt->get_path(); + RAFT_EXPECTS(!graph_path.empty(), "Unable to get path from graph file descriptor"); + RAFT_EXPECTS(!mapping_path.empty(), "Unable to get path from mapping file descriptor"); + + auto graph_npy = open_npy_file(graph_path); + auto mapping_npy = open_npy_file(mapping_path); + + validate_npy_file(graph_npy, graph_path, "Graph file"); + validate_npy_file(mapping_npy, mapping_path, "Mapping file"); + RAFT_EXPECTS(graph_npy.shape.size() == 2, "Graph file should be 2D"); + RAFT_EXPECTS(mapping_npy.shape.size() == 1, "Mapping file should be 1D"); + const auto n_rows = graph_npy.shape[0]; + const auto degree = graph_npy.shape[1]; + RAFT_EXPECTS(mapping_npy.shape[0] == n_rows, + "Mapping elements (%zu) != graph rows (%zu)", + mapping_npy.shape[0], + n_rows); + RAFT_EXPECTS(n_rows == static_cast(index_.size()), + "Graph rows (%zu) != index size (%zu)", + n_rows, + static_cast(index_.size())); + RAFT_EXPECTS(degree == static_cast(index_.graph_degree()), + "Graph degree (%zu) != index graph degree (%zu)", + degree, + static_cast(index_.graph_degree())); + RAFT_EXPECTS(degree > 0, "Graph degree must be nonzero"); + RAFT_EXPECTS( + degree <= maxM0, "Base graph degree (%zu) must not exceed HNSW maxM0 (%zu)", degree, maxM0); + RAFT_EXPECTS(base_link_row_bytes >= sizeof(hnswlib::linklistsizeint) + maxM0 * sizeof(IdxT), + "Base link row size is too small"); + + RAFT_LOG_INFO("Layered HNSW artifact: loading ACE row mapping (%zu rows)", n_rows); + const auto mapping_start_time = std::chrono::steady_clock::now(); + const auto mapping_bytes = n_rows * sizeof(IdxT); + std::vector reordered_to_original(n_rows); + cuvs::util::read_large_file( + mapping_npy.fd, reordered_to_original.data(), mapping_bytes, mapping_npy.header_size); + for (size_t reordered_id = 0; reordered_id < n_rows; ++reordered_id) { + const auto original_id = static_cast(reordered_to_original[reordered_id]); + RAFT_EXPECTS(original_id < n_rows, + "Invalid original id %zu in ACE dataset mapping at row %zu", + original_id, + reordered_id); + } + const auto mapping_elapsed_ms = elapsed_ms_since(mapping_start_time); + RAFT_LOG_INFO("Layered HNSW artifact: ACE row mapping loaded in %ld ms (%.2f GiB, %.2f GiB/s)", + mapping_elapsed_ms, + to_gib(mapping_bytes), + throughput_gib_per_s(mapping_bytes, mapping_elapsed_ms)); + + RAFT_LOG_INFO( + "Layered HNSW artifact: writing base topology with source-sequential graph reads " + "(rows=%zu degree=%zu maxM0=%zu)", + n_rows, + degree, + maxM0); + + const auto total_start_time = std::chrono::steady_clock::now(); + const auto graph_row_bytes = degree * sizeof(IdxT); + const auto base_node_row_bytes = sizeof(IdxT); + const auto base_topology_row_size = graph_row_bytes + base_node_row_bytes + base_link_row_bytes; + const size_t target_batch_bytes = 64 * 1024 * 1024; + const size_t batch_size = std::max(1, target_batch_bytes / base_topology_row_size); + auto graph_buffer = raft::make_host_matrix(static_cast(batch_size), + static_cast(degree)); + std::vector base_node_buffer(batch_size); + std::vector base_link_buffer(batch_size * base_link_row_bytes); + + size_t graph_bytes_read = 0; + size_t node_bytes_written = 0; + size_t link_bytes_written = 0; + RAFT_LOG_INFO( + "Layered HNSW artifact: base topology batch_size=%zu graph_row=%zu bytes node_row=%zu bytes " + "link_row=%zu bytes", + batch_size, + graph_row_bytes, + base_node_row_bytes, + base_link_row_bytes); + size_t next_report_percent = 10; + for (size_t source_start = 0; source_start < n_rows; source_start += batch_size) { + const auto current_batch_size = std::min(batch_size, n_rows - source_start); + const auto batch_bytes = current_batch_size * graph_row_bytes; + cuvs::util::read_large_file(graph_npy.fd, + graph_buffer.data_handle(), + batch_bytes, + graph_npy.header_size + source_start * graph_row_bytes); + graph_bytes_read += batch_bytes; + + // Rows stay in source (ACE-reordered) order so the artifact is written sequentially on the + // build node. Each row records its original ID in base_nodes and remaps its neighbors to + // original IDs. The search node scatters each row into get_linklist0(original_id) on load. + bool invalid_neighbor = false; +#pragma omp parallel for reduction(|| : invalid_neighbor) + for (int64_t batch_idx = 0; batch_idx < static_cast(current_batch_size); ++batch_idx) { + const auto source_row = source_start + static_cast(batch_idx); + const auto original_id = static_cast(reordered_to_original[source_row]); + base_node_buffer[batch_idx] = static_cast(original_id); + auto* dst_row_ptr = base_link_buffer.data() + batch_idx * base_link_row_bytes; + hnswlib::linklistsizeint list_count = static_cast(degree); + std::memcpy(dst_row_ptr, &list_count, sizeof(list_count)); + auto* dst = reinterpret_cast(dst_row_ptr + sizeof(hnswlib::linklistsizeint)); + auto* src = graph_buffer.data_handle() + batch_idx * degree; + for (size_t j = 0; j < degree; ++j) { + const auto neighbor = static_cast(src[j]); + if (neighbor >= n_rows) { + invalid_neighbor = true; + continue; + } + dst[j] = reordered_to_original[neighbor]; + } + const auto written_bytes = sizeof(hnswlib::linklistsizeint) + degree * sizeof(IdxT); + if (written_bytes < base_link_row_bytes) { + std::memset(dst_row_ptr + written_bytes, 0, base_link_row_bytes - written_bytes); + } + } + RAFT_EXPECTS(!invalid_neighbor, "Invalid reordered neighbor id in ACE graph"); + + const auto current_node_bytes = current_batch_size * sizeof(IdxT); + const auto current_link_bytes = current_batch_size * base_link_row_bytes; + cuvs::util::write_large_file(output_fd, + base_node_buffer.data(), + current_node_bytes, + base_nodes_offset + source_start * sizeof(IdxT)); + cuvs::util::write_large_file(output_fd, + base_link_buffer.data(), + current_link_bytes, + base_links_offset + source_start * base_link_row_bytes); + node_bytes_written += current_node_bytes; + link_bytes_written += current_link_bytes; + + const auto completed_rows = source_start + current_batch_size; + const auto report_percent = progress_step_10(completed_rows, n_rows); + if (report_percent >= next_report_percent) { + const auto elapsed = elapsed_ms_since(total_start_time); + RAFT_LOG_INFO( + "Layered HNSW artifact: base topology progress %zu%% (%zu/%zu rows, read=%.2f GiB " + "write=%.2f GiB, %ld ms)", + report_percent, + completed_rows, + n_rows, + to_gib(graph_bytes_read), + to_gib(node_bytes_written + link_bytes_written), + elapsed); + next_report_percent = report_percent + 10; + } + } + + const auto total_elapsed_ms = elapsed_ms_since(total_start_time); + RAFT_LOG_INFO( + "Layered HNSW artifact: base topology written in %ld ms (graph_read=%.2f GiB %.2f GiB/s, " + "artifact_write=%.2f GiB %.2f GiB/s, nodes=%.2f GiB links=%.2f GiB)", + total_elapsed_ms, + to_gib(graph_bytes_read), + throughput_gib_per_s(graph_bytes_read, total_elapsed_ms), + to_gib(node_bytes_written + link_bytes_written), + throughput_gib_per_s(node_bytes_written + link_bytes_written, total_elapsed_ms), + to_gib(node_bytes_written), + to_gib(link_bytes_written)); +} + +template +auto serialize_to_layered_hnswlib_from_disk( + raft::resources const& res, + const cuvs::neighbors::hnsw::index_params& params, + const CagraIndexT& index_, + raft::host_matrix_view dataset) -> std::string +{ + using IdxT = typename CagraIndexT::index_type; + static_assert(is_cagra_hnsw_export_index_v, + "Layered HNSW serialization expects a CAGRA index type"); + raft::common::nvtx::range fun_scope("hnsw::serialize_layered"); + const auto total_start_time = std::chrono::steady_clock::now(); + + RAFT_EXPECTS(index_.graph_fd().has_value() && index_.mapping_fd().has_value(), + "Layered HNSW serialization requires a disk-backed ACE graph and mapping."); + RAFT_EXPECTS(static_cast(dataset.extent(0)) == static_cast(index_.size()), + "Dataset rows (%zu) must match index size (%zu)", + static_cast(dataset.extent(0)), + static_cast(index_.size())); + RAFT_EXPECTS(static_cast(dataset.extent(1)) == static_cast(index_.dim()), + "Dataset cols (%zu) must match index dimensions (%zu)", + static_cast(dataset.extent(1)), + static_cast(index_.dim())); + + const auto graph_path = index_.graph_fd()->get_path(); + RAFT_EXPECTS(!graph_path.empty(), "Unable to get path from graph file descriptor"); + const auto ace_dir = std::filesystem::path(graph_path).parent_path(); + const auto artifact_file = ace_dir / "hnsw_index.cuvs"; + std::filesystem::create_directories(ace_dir); + + auto n_rows = static_cast(index_.size()); + auto dim = static_cast(index_.dim()); + auto graph_degree_int = static_cast(index_.graph_degree()); + + auto hnsw_index = + std::make_unique>(index_.dim(), index_.metric(), params.hierarchy); + auto appr_algo = std::make_unique::type>>( + hnsw_index->get_space(), 1, (graph_degree_int + 1) / 2, params.ef_construction); + + RAFT_LOG_INFO("Layered HNSW artifact: generating hierarchy levels"); + const auto hierarchy_start_time = std::chrono::steady_clock::now(); + auto hierarchy = make_random_hnsw_level_plan(res, n_rows, *appr_algo, nullptr); + const auto hierarchy_elapsed_ms = elapsed_ms_since(hierarchy_start_time); + RAFT_LOG_INFO( + "Layered HNSW artifact: hierarchy levels generated in %ld ms (max_level=%d promoted=%zu)", + hierarchy_elapsed_ms, + hierarchy.max_level(), + hierarchy.promoted_count()); + + layered_hnsw_file_metadata metadata; + metadata.n_rows = n_rows; + metadata.dim = dim; + metadata.M = appr_algo->M_; + metadata.maxM = appr_algo->maxM_; + metadata.maxM0 = appr_algo->maxM0_; + metadata.ef_construction = appr_algo->ef_construction_; + metadata.mult = appr_algo->mult_; + metadata.maxlevel = hierarchy.max_level(); + metadata.enterpoint_node = + static_cast(hierarchy.row_ids_by_level(hierarchy.row_ids_by_level.size() - 1)); + metadata.base_degree = static_cast(graph_degree_int); + metadata.levels_bytes = n_rows * sizeof(uint8_t); + metadata.base_nodes_bytes = n_rows * sizeof(IdxT); + metadata.base_link_row_bytes = appr_algo->size_links_level0_; + metadata.base_links_bytes = n_rows * metadata.base_link_row_bytes; + metadata.upper_link_row_bytes = appr_algo->size_links_per_element_; + size_t next_upper_node_offset = 0; + size_t next_upper_link_offset = 0; + for (size_t pt_level = 1; pt_level < hierarchy.hist.size(); pt_level++) { + auto start_idx = hierarchy.offsets[pt_level - 1]; + auto row_count = n_rows - start_idx; + auto layer_degree = hnsw_upper_layer_degree(row_count, appr_algo->M_); + metadata.layers.push_back( + {pt_level, row_count, layer_degree, next_upper_node_offset, next_upper_link_offset}); + next_upper_node_offset += row_count; + next_upper_link_offset += row_count; + } + metadata.upper_nodes_count = next_upper_node_offset; + metadata.upper_nodes_bytes = metadata.upper_nodes_count * sizeof(IdxT); + metadata.upper_links_bytes = next_upper_link_offset * metadata.upper_link_row_bytes; + + const auto header = make_layered_hnsw_header(metadata, index_.metric()); + std::vector layer_descriptors(metadata.layers.size()); + for (size_t i = 0; i < metadata.layers.size(); ++i) { + const auto& layer = metadata.layers[i]; + layer_descriptors[i] = {static_cast(layer.level), + static_cast(layer.row_count), + static_cast(layer.degree), + static_cast(layer.node_offset), + static_cast(layer.link_offset)}; + } + const auto descriptors_offset = sizeof(layered_hnsw_file_header); + const auto descriptors_bytes = layer_descriptors.size() * sizeof(layered_hnsw_layer_descriptor); + + const auto payload_offset = + align_up(descriptors_offset + descriptors_bytes, layered_hnsw_alignment); + const auto levels_offset = payload_offset; + const auto base_nodes_offset = levels_offset + metadata.levels_bytes; + const auto base_links_offset = base_nodes_offset + metadata.base_nodes_bytes; + const auto upper_nodes_offset = base_links_offset + metadata.base_links_bytes; + const auto upper_links_offset = upper_nodes_offset + metadata.upper_nodes_bytes; + const auto final_file_size = upper_links_offset + metadata.upper_links_bytes; + + exclusive_hnsw_temp_file output(artifact_file); + cuvs::util::file_descriptor artifact_fd(output.temporary_path(), O_RDWR); + const auto fallocate_result = posix_fallocate(artifact_fd.get(), 0, final_file_size); + RAFT_EXPECTS(fallocate_result == 0, + "Failed to pre-allocate layered HNSW artifact %s: %s", + artifact_file.string().c_str(), + std::strerror(fallocate_result)); + cuvs::util::write_large_file(artifact_fd, &header, sizeof(header), 0); + if (descriptors_bytes > 0) { + cuvs::util::write_large_file( + artifact_fd, layer_descriptors.data(), descriptors_bytes, descriptors_offset); + } + + const auto levels_start_time = std::chrono::steady_clock::now(); + cuvs::util::write_large_file( + artifact_fd, hierarchy.levels.data_handle(), metadata.levels_bytes, levels_offset); + const auto levels_elapsed_ms = elapsed_ms_since(levels_start_time); + RAFT_LOG_INFO("Layered HNSW artifact: levels section written in %ld ms (%.2f GiB, %.2f GiB/s)", + levels_elapsed_ms, + to_gib(metadata.levels_bytes), + throughput_gib_per_s(metadata.levels_bytes, levels_elapsed_ms)); + + RAFT_LOG_INFO("Layered HNSW artifact: writing hnswlib-ready base topology section"); + const auto layer0_start_time = std::chrono::steady_clock::now(); + write_layered_base_links_from_disk(index_, + artifact_fd, + base_nodes_offset, + base_links_offset, + metadata.base_link_row_bytes, + metadata.maxM0); + const auto layer0_elapsed_ms = elapsed_ms_since(layer0_start_time); + static_cast(layer0_elapsed_ms); + RAFT_LOG_INFO( + "Layered HNSW artifact: base topology section written in %ld ms (%.2f GiB, %.2f GiB/s)", + layer0_elapsed_ms, + to_gib(metadata.base_nodes_bytes + metadata.base_links_bytes), + throughput_gib_per_s(metadata.base_nodes_bytes + metadata.base_links_bytes, layer0_elapsed_ms)); + + size_t upper_graph_bytes_written = 0; + if (hierarchy.hist.size() > 1) { + RAFT_LOG_INFO("Layered HNSW artifact: gathering promoted vectors"); + const auto gather_start_time = std::chrono::steady_clock::now(); + auto host_query_set = + raft::make_host_matrix(static_cast(hierarchy.promoted_count()), dim); +#pragma omp parallel for + for (int64_t i = 0; i < static_cast(n_rows); i++) { + if (hierarchy.levels(i) > 0) { + const auto query_row = hierarchy.position_by_row_id(i) - hierarchy.hist[0]; + auto* dst = host_query_set.data_handle() + query_row * dim; + std::copy(&dataset(i, 0), &dataset(i, 0) + dim, dst); + } + } + const auto gather_elapsed_ms = elapsed_ms_since(gather_start_time); + const auto gathered_bytes = hierarchy.promoted_count() * dim * sizeof(T); + RAFT_LOG_INFO( + "Layered HNSW artifact: promoted vectors gathered in %ld ms (%.2f GiB copied, %.2f GiB/s)", + gather_elapsed_ms, + to_gib(gathered_bytes), + throughput_gib_per_s(gathered_bytes, gather_elapsed_ms)); + + auto promoted_dataset = raft::make_host_matrix_view( + host_query_set.data_handle(), host_query_set.extent(0), host_query_set.extent(1)); + const auto upper_layers_start_time = std::chrono::steady_clock::now(); + build_hnsw_upper_layer_graphs( + res, + promoted_dataset, + hierarchy, + appr_algo->M_, + index_.metric(), + [&](size_t pt_level, size_t start_idx, auto& host_neighbors) { + const auto& layer = metadata.layers[pt_level - 1]; + RAFT_LOG_DEBUG("Layered HNSW artifact: writing upper layer level=%zu rows=%zu degree=%zu", + layer.level, + layer.row_count, + layer.degree); + const auto layer_write_start_time = std::chrono::steady_clock::now(); + const size_t target_batch_bytes = 64 * 1024 * 1024; + const size_t row_bytes = sizeof(IdxT) + metadata.upper_link_row_bytes; + const size_t batch_size = std::max(1, target_batch_bytes / row_bytes); + std::vector node_buffer(batch_size); + std::vector link_buffer(batch_size * metadata.upper_link_row_bytes); + for (size_t batch_start = 0; batch_start < layer.row_count; batch_start += batch_size) { + const auto current_batch_size = std::min(batch_size, layer.row_count - batch_start); + std::fill(link_buffer.begin(), + link_buffer.begin() + current_batch_size * metadata.upper_link_row_bytes, + 0); + // Upper-layer rows stay in promoted (per-level) order for sequential writes. node_buffer + // records each node's original ID and neighbors are remapped to original IDs. The search + // node scatters each row into get_linklist(original_id, level) on load. +#pragma omp parallel for + for (int64_t batch_idx = 0; batch_idx < static_cast(current_batch_size); + ++batch_idx) { + const auto row = batch_start + static_cast(batch_idx); + node_buffer[batch_idx] = static_cast(hierarchy.row_ids_by_level(start_idx + row)); + auto* link_row = link_buffer.data() + batch_idx * metadata.upper_link_row_bytes; + hnswlib::linklistsizeint list_count = + static_cast(layer.degree); + std::memcpy(link_row, &list_count, sizeof(list_count)); + auto* dst = reinterpret_cast(link_row + sizeof(hnswlib::linklistsizeint)); + if (layer.degree > 0) { + auto* src = host_neighbors.data_handle() + row * layer.degree; + for (size_t j = 0; j < layer.degree; ++j) { + dst[j] = static_cast(hierarchy.row_ids_by_level(src[j] + start_idx)); + } + } + } + cuvs::util::write_large_file( + artifact_fd, + node_buffer.data(), + current_batch_size * sizeof(IdxT), + upper_nodes_offset + (layer.node_offset + batch_start) * sizeof(IdxT)); + cuvs::util::write_large_file( + artifact_fd, + link_buffer.data(), + current_batch_size * metadata.upper_link_row_bytes, + upper_links_offset + (layer.link_offset + batch_start) * metadata.upper_link_row_bytes); + } + const auto layer_bytes = layer.row_count * (sizeof(IdxT) + metadata.upper_link_row_bytes); + upper_graph_bytes_written += layer_bytes; + const auto layer_write_elapsed_ms = elapsed_ms_since(layer_write_start_time); + static_cast(layer_write_elapsed_ms); + RAFT_LOG_INFO( + "Layered HNSW artifact: upper layer level=%zu written in %ld ms (%.2f GiB, %.2f GiB/s)", + layer.level, + layer_write_elapsed_ms, + to_gib(layer_bytes), + throughput_gib_per_s(layer_bytes, layer_write_elapsed_ms)); + }); + const auto upper_layers_elapsed_ms = elapsed_ms_since(upper_layers_start_time); + RAFT_LOG_INFO( + "Layered HNSW artifact: upper layers generated and written in %ld ms (%.2f GiB, %.2f " + "GiB/s)", + upper_layers_elapsed_ms, + to_gib(upper_graph_bytes_written), + throughput_gib_per_s(upper_graph_bytes_written, upper_layers_elapsed_ms)); + } + + const auto total_elapsed_ms = elapsed_ms_since(total_start_time); + RAFT_LOG_INFO("Layered HNSW artifact: wrote %s in %ld ms (artifact=%.2f GiB, %.2f GiB/s)", + artifact_file.string().c_str(), + total_elapsed_ms, + to_gib(final_file_size), + throughput_gib_per_s(final_file_size, total_elapsed_ms)); + artifact_fd.close(); + output.publish(); + return artifact_file.string(); +} + // Source-agnostic core that streams a CAGRA index into hnswlib format on disk. // The disk-backed and in-memory variants differ only in the `read_batch` callable, // which fills the provided host buffers (graph, dataset, labels) for a given row range. @@ -522,43 +1380,21 @@ void serialize_to_hnswlib_batched(raft::resources const& res, bool create_hierarchy = params.hierarchy != HnswHierarchy::NONE; - // create hierarchy order - // sort the points by levels - // roll dice & build histogram - std::vector hist; - std::vector order(n_rows); - std::vector order_bw(n_rows); - std::vector levels(n_rows); - std::vector offsets; - + hnsw_level_plan hierarchy{res}; if (create_hierarchy) { RAFT_LOG_INFO("Sort points by levels"); - for (int64_t i = 0; i < n_rows; i++) { - auto pt_level = appr_algo->getRandomLevel(appr_algo->mult_); - while (pt_level >= static_cast(hist.size())) - hist.push_back(0); - hist[pt_level]++; - levels[i] = pt_level; - } - - // accumulate - offsets.resize(hist.size() + 1, 0); - for (size_t i = 0; i < hist.size() - 1; i++) { - offsets[i + 1] = offsets[i] + hist[i]; - RAFT_LOG_INFO("Level %zu : %zu", i + 1, size_t(n_rows) - offsets[i + 1]); - } - - // fw/bw indices - for (int64_t i = 0; i < n_rows; i++) { - auto pt_level = levels[i]; - order_bw[i] = offsets[pt_level]; - order[offsets[pt_level]++] = i; - } + hierarchy = make_random_hnsw_level_plan(res, n_rows, *appr_algo, "Level "); } + const auto& hist = hierarchy.hist; + const auto& row_ids_by_level = hierarchy.row_ids_by_level; + const auto& position_by_row_id = hierarchy.position_by_row_id; + const auto& levels = hierarchy.levels; + const auto& offsets = hierarchy.offsets; // set last point of the highest level as the entry point - appr_algo->enterpoint_node_ = create_hierarchy ? order.back() : n_rows / 2; - appr_algo->maxlevel_ = create_hierarchy ? hist.size() - 1 : 1; + appr_algo->enterpoint_node_ = create_hierarchy ? row_ids_by_level(row_ids_by_level.size() - 1) + : static_cast(n_rows / 2); + appr_algo->maxlevel_ = create_hierarchy ? hierarchy.max_level() : 1; // write header information RAFT_LOG_DEBUG("Writing HNSW header: offsetLevel0=%zu, n_rows=%zu, size_data_per_element=%zu", @@ -650,11 +1486,11 @@ void serialize_to_hnswlib_batched(raft::resources const& res, const T* data_row = &dataset_buffer(batch_idx, 0); os.write(reinterpret_cast(data_row), sizeof(T) * dim); - if (create_hierarchy && levels[i] > 0) { - // position in query: order_bw[i]-hist[0] + if (create_hierarchy && levels(i) > 0) { + // position in query: position_by_row_id[i]-hist[0] std::copy(data_row, data_row + dim, - reinterpret_cast(&host_query_set(order_bw[i] - hist[0], 0))); + reinterpret_cast(&host_query_set(position_by_row_id(i) - hist[0], 0))); } // assign original label @@ -717,7 +1553,7 @@ void serialize_to_hnswlib_batched(raft::resources const& res, start_clock = std::chrono::system_clock::now(); for (int64_t i = 0; i < n_rows; i++) { - size_t cur_level = create_hierarchy ? levels[i] : 0; + size_t cur_level = create_hierarchy ? levels(i) : 0; unsigned int linkListSize = create_hierarchy && cur_level > 0 ? appr_algo->size_links_per_element_ * cur_level : 0; os.write(reinterpret_cast(&linkListSize), sizeof(int)); @@ -725,13 +1561,13 @@ void serialize_to_hnswlib_batched(raft::resources const& res, if (linkListSize) { for (size_t pt_level = 1; pt_level <= cur_level; pt_level++) { auto neighbor_view = host_neighbors[pt_level - 1].view(); - auto my_row = order_bw[i] - offsets[pt_level - 1]; + auto my_row = position_by_row_id(i) - offsets[pt_level - 1]; IdxT* neighbors = &neighbor_view(my_row, 0); unsigned int extent = neighbor_view.extent(1); os.write(reinterpret_cast(&extent), sizeof(int)); for (unsigned int j = 0; j < extent; j++) { - const IdxT converted = order[neighbors[j] + offsets[pt_level - 1]]; + const IdxT converted = row_ids_by_level(neighbors[j] + offsets[pt_level - 1]); os.write(reinterpret_cast(&converted), sizeof(IdxT)); } auto remainder = appr_algo->M_ - neighbor_view.extent(1); @@ -1368,6 +2204,18 @@ inline std::pair get_available_memory( return std::make_pair(available_host_memory, available_device_memory); } +inline void validate_output_format(const index_params& params) +{ + switch (params.output_format) { + case HnswOutputFormat::HNSWLIB: return; + case HnswOutputFormat::GRAPH_ONLY: + RAFT_EXPECTS(params.hierarchy == HnswHierarchy::GPU, + "GRAPH_ONLY requires HnswHierarchy::GPU"); + return; + } + RAFT_FAIL("Unsupported HNSW output format"); +} + template requires is_cagra_hnsw_export_index_v std::unique_ptr> from_cagra( @@ -1376,6 +2224,8 @@ std::unique_ptr> from_cagra( CagraIndexT const& cagra_index, std::optional> dataset) { + validate_output_format(params); + if constexpr (is_host_cagra_hnsw_export_index_v) { if (!cagra_index.dataset_fd().has_value() && !dataset.has_value()) { RAFT_FAIL("hnsw::from_cagra requires dataset for host CAGRA index"); @@ -1397,6 +2247,18 @@ std::unique_ptr> from_cagra( std::filesystem::exists(index_directory) && std::filesystem::is_directory(index_directory), "Directory '%s' does not exist", index_directory.c_str()); + if (params.output_format == HnswOutputFormat::GRAPH_ONLY) { + RAFT_EXPECTS(dataset.has_value(), + "Layered HNSW serialization requires the original-order dataset."); + auto artifact_path = + serialize_to_layered_hnswlib_from_disk(res, params, cagra_index, dataset.value()); + + auto hnsw_index = std::make_unique>( + cagra_index.dim(), cagra_index.metric(), params.hierarchy, params.output_format); + hnsw_index->set_file_descriptor(cuvs::util::file_descriptor(artifact_path, O_RDONLY)); + return hnsw_index; + } + std::string index_filename = (std::filesystem::path(index_directory) / "hnsw_index.bin").string(); exclusive_hnsw_output_file output(index_filename); @@ -1405,8 +2267,8 @@ std::unique_ptr> from_cagra( output.publish(); // Create an empty HNSW index that holds the file descriptor - auto hnsw_index = - std::make_unique>(cagra_index.dim(), cagra_index.metric(), params.hierarchy); + auto hnsw_index = std::make_unique>( + cagra_index.dim(), cagra_index.metric(), params.hierarchy, params.output_format); // Open file descriptor for the HNSW index file and transfer ownership to the index hnsw_index->set_file_descriptor(cuvs::util::file_descriptor(index_filename, O_RDONLY)); @@ -1416,11 +2278,15 @@ std::unique_ptr> from_cagra( return hnsw_index; } + RAFT_EXPECTS(params.output_format == HnswOutputFormat::HNSWLIB, + "GRAPH_ONLY requires disk-backed ACE build artifacts"); + // In-memory CAGRA index: the resulting HNSW index might still not fit in host memory. // Estimate its host footprint and, if it does not fit, spill it to disk via // serialize_to_hnswlib_from_inmem instead of constructing it in RAM (NONE/GPU only; - // the CPU hierarchy is not supported by the batched serializer). - if (params.hierarchy != HnswHierarchy::CPU) { + // the CPU hierarchy is not supported by the batched serializer, and GRAPH_ONLY is only + // produced from disk-backed ACE artifacts handled above). + if (params.hierarchy == HnswHierarchy::NONE || params.hierarchy == HnswHierarchy::GPU) { int64_t n_rows = dataset.has_value() ? dataset->extent(0) : cagra_index.size(); int64_t dim = dataset.has_value() ? dataset->extent(1) : cagra_index.dim(); int graph_degree_int = static_cast(cagra_index.graph().extent(1)); @@ -1626,6 +2492,13 @@ void serialize(raft::resources const& res, const std::string& filename, const in "Disk-based index file does not exist: %s", source_path.c_str()); + if (idx_impl->output_format() == HnswOutputFormat::GRAPH_ONLY) { + copy_file_overwrite(source_path, filename); + RAFT_LOG_INFO( + "Copied layered HNSW index from %s to %s", source_path.c_str(), filename.c_str()); + return; + } + // Copy the file to the new location std::filesystem::copy_file( source_path, filename, std::filesystem::copy_options::overwrite_existing); @@ -1639,6 +2512,378 @@ void serialize(raft::resources const& res, const std::string& filename, const in hnswlib_index->saveIndex(filename); } +template +auto deserialize_layered_hnswlib(raft::resources const& res, + const std::string& artifact_path, + const std::string& dataset_path) -> std::unique_ptr> +{ + common::nvtx::range fun_scope("hnsw::deserialize_layered"); + const auto total_start_time = std::chrono::steady_clock::now(); + const auto metadata_start_time = std::chrono::steady_clock::now(); + cuvs::util::file_descriptor artifact_fd(artifact_path, O_RDONLY); + layered_hnsw_file_header header{}; + cuvs::util::read_large_file(artifact_fd, &header, sizeof(header), 0); + RAFT_EXPECTS(std::strncmp(header.magic, layered_hnsw_magic, sizeof(header.magic)) == 0, + "Invalid layered HNSW artifact magic: %s", + artifact_path.c_str()); + RAFT_EXPECTS(header.version == layered_hnsw_version, + "Unsupported layered HNSW artifact version %u", + header.version); + const auto construction_dtype = static_cast(header.construction_dtype); + RAFT_EXPECTS(is_supported_layered_dtype(construction_dtype), + "Layered HNSW artifact contains unsupported construction dtype code %u", + header.construction_dtype); + const auto metric = static_cast(header.metric); + RAFT_EXPECTS(metric == cuvs::distance::DistanceType::L2Expanded || + metric == cuvs::distance::DistanceType::InnerProduct, + "Layered HNSW artifact contains unsupported metric code %u", + header.metric); + RAFT_EXPECTS(header.dim <= static_cast(std::numeric_limits::max()), + "Layered HNSW artifact dimension (%zu) exceeds supported maximum (%d)", + static_cast(header.dim), + std::numeric_limits::max()); + const auto dim = static_cast(header.dim); + + const auto artifact_size = static_cast(std::filesystem::file_size(artifact_path)); + const auto descriptors_offset = sizeof(layered_hnsw_file_header); + const auto descriptors_bytes = + static_cast(header.num_layers) * sizeof(layered_hnsw_layer_descriptor); + RAFT_EXPECTS(descriptors_offset + descriptors_bytes <= artifact_size, + "Layered HNSW layer descriptors are outside artifact: offset=%zu size=%zu " + "artifact=%zu", + descriptors_offset, + descriptors_bytes, + artifact_size); + + std::vector layer_descriptors(header.num_layers); + if (descriptors_bytes > 0) { + cuvs::util::read_large_file( + artifact_fd, layer_descriptors.data(), descriptors_bytes, descriptors_offset); + } + auto metadata = layered_hnsw_metadata_from_header(header); + metadata.layers.reserve(layer_descriptors.size()); + for (const auto& descriptor : layer_descriptors) { + metadata.layers.push_back({static_cast(descriptor.level), + static_cast(descriptor.row_count), + static_cast(descriptor.degree), + static_cast(descriptor.node_offset), + static_cast(descriptor.link_offset)}); + } + const auto metadata_elapsed_ms = elapsed_ms_since(metadata_start_time); + + RAFT_EXPECTS(metadata.n_rows > 0, "Layered HNSW artifact must contain at least one row"); + RAFT_EXPECTS(metadata.dim > 0, "Layered HNSW artifact must contain at least one dimension"); + RAFT_EXPECTS(metadata.enterpoint_node >= 0 && + static_cast(metadata.enterpoint_node) < metadata.n_rows, + "Layered HNSW enterpoint node (%d) outside valid range [0, %zu)", + metadata.enterpoint_node, + metadata.n_rows); + RAFT_EXPECTS(metadata.layers.size() == static_cast(metadata.maxlevel), + "Layered HNSW artifact has %zu upper layers, expected %d", + metadata.layers.size(), + metadata.maxlevel); + + RAFT_EXPECTS(!dataset_path.empty(), "Layered HNSW deserialization requires a dataset filename"); + + const auto payload_offset = + align_up(descriptors_offset + descriptors_bytes, layered_hnsw_alignment); + const auto levels_offset = payload_offset; + const auto base_nodes_offset = levels_offset + metadata.levels_bytes; + const auto base_links_offset = base_nodes_offset + metadata.base_nodes_bytes; + const auto upper_nodes_offset = base_links_offset + metadata.base_links_bytes; + const auto upper_links_offset = upper_nodes_offset + metadata.upper_nodes_bytes; + const auto expected_file_size = upper_links_offset + metadata.upper_links_bytes; + RAFT_EXPECTS(artifact_size >= expected_file_size, + "Layered HNSW artifact is truncated: expected at least %zu bytes, got %zu", + expected_file_size, + artifact_size); + RAFT_LOG_INFO("Layered HNSW load: metadata read in %ld ms (rows=%zu dim=%zu artifact=%.2f GiB)", + metadata_elapsed_ms, + metadata.n_rows, + metadata.dim, + to_gib(artifact_size)); + + const auto dataset_open_start_time = std::chrono::steady_clock::now(); + auto dataset_file = open_layered_dataset_file(dataset_path); + const auto dataset_open_elapsed_ms = elapsed_ms_since(dataset_open_start_time); + + RAFT_EXPECTS(dataset_file.shape.size() == 2 && dataset_file.shape[0] == metadata.n_rows && + dataset_file.shape[1] == metadata.dim, + "Layered HNSW dataset shape mismatch: artifact rows=%zu dim=%zu, dataset rows=%zu " + "dim=%zu path=%s", + metadata.n_rows, + metadata.dim, + dataset_file.shape.size() > 0 ? dataset_file.shape[0] : 0, + dataset_file.shape.size() > 1 ? dataset_file.shape[1] : 0, + dataset_path.c_str()); + validate_npy_file(dataset_file, dataset_path, "Layered HNSW dataset"); + RAFT_EXPECTS(metadata.levels_bytes == metadata.n_rows * sizeof(uint8_t), + "Layered HNSW levels section size mismatch"); + RAFT_EXPECTS(metadata.base_nodes_bytes == metadata.n_rows * sizeof(uint32_t), + "Layered HNSW base node section size mismatch"); + RAFT_EXPECTS(metadata.base_links_bytes == metadata.n_rows * metadata.base_link_row_bytes, + "Layered HNSW base links section size mismatch"); + RAFT_EXPECTS(metadata.upper_nodes_bytes == metadata.upper_nodes_count * sizeof(uint32_t), + "Layered HNSW upper node section size mismatch"); + RAFT_EXPECTS( + metadata.upper_links_bytes == metadata.upper_nodes_count * metadata.upper_link_row_bytes, + "Layered HNSW upper link section size mismatch"); + RAFT_EXPECTS(metadata.base_degree <= metadata.maxM0, + "Layered HNSW base degree (%zu) exceeds maxM0 (%zu)", + metadata.base_degree, + metadata.maxM0); + + RAFT_LOG_INFO("Layered HNSW load: dataset header validated in %ld ms (%s)", + dataset_open_elapsed_ms, + dataset_path.c_str()); + + const auto dataset_total_bytes = metadata.n_rows * metadata.dim * sizeof(T); + const auto deserialize_progress_total_bytes = + metadata.levels_bytes + dataset_total_bytes + metadata.base_nodes_bytes + + metadata.base_links_bytes + metadata.upper_nodes_bytes + metadata.upper_links_bytes; + size_t deserialize_progress_bytes = 0; + size_t deserialize_next_progress_percent = 10; + auto log_deserialize_progress = [&](size_t bytes_loaded) { + if (deserialize_progress_total_bytes == 0) { return; } + deserialize_progress_bytes = + std::min(deserialize_progress_total_bytes, deserialize_progress_bytes + bytes_loaded); + const auto progress_percent = + (deserialize_progress_bytes * 100) / deserialize_progress_total_bytes; + const auto progress_step = std::min(100, (progress_percent / 10) * 10); + if (progress_step >= deserialize_next_progress_percent) { + RAFT_LOG_INFO("Layered HNSW load: deserialize progress %zu%% (%.2f/%.2f GiB)", + progress_step, + to_gib(deserialize_progress_bytes), + to_gib(deserialize_progress_total_bytes)); + deserialize_next_progress_percent = progress_step + 10; + } + }; + + const auto levels_start_time = std::chrono::steady_clock::now(); + std::vector levels_u8(metadata.n_rows); + cuvs::util::read_large_file(artifact_fd, levels_u8.data(), metadata.levels_bytes, levels_offset); + log_deserialize_progress(metadata.levels_bytes); + const auto max_level_in_levels = *std::max_element(levels_u8.begin(), levels_u8.end()); + RAFT_EXPECTS(static_cast(max_level_in_levels) == metadata.maxlevel, + "Layered HNSW levels max level (%d) does not match artifact maxlevel (%d)", + static_cast(max_level_in_levels), + metadata.maxlevel); + const auto levels_elapsed_ms = elapsed_ms_since(levels_start_time); + RAFT_LOG_INFO("Layered HNSW load: levels read in %ld ms (%.2f MiB)", + levels_elapsed_ms, + static_cast(metadata.levels_bytes) / (1024.0 * 1024.0)); + + const auto allocation_start_time = std::chrono::steady_clock::now(); + auto hnsw_index = + std::make_unique>(dim, metric, HnswHierarchy::GPU, HnswOutputFormat::GRAPH_ONLY); + auto appr_algo = std::make_unique::type>>( + hnsw_index->get_space(), metadata.n_rows, metadata.M, metadata.ef_construction); + appr_algo->cur_element_count = metadata.n_rows; + appr_algo->maxlevel_ = metadata.maxlevel; + appr_algo->enterpoint_node_ = metadata.enterpoint_node; + std::fill_n(appr_algo->linkLists_, metadata.n_rows, nullptr); + RAFT_EXPECTS(appr_algo->size_links_level0_ == metadata.base_link_row_bytes, + "Layered HNSW base link row size mismatch"); + RAFT_EXPECTS(appr_algo->size_links_per_element_ == metadata.upper_link_row_bytes, + "Layered HNSW upper link row size mismatch"); + RAFT_EXPECTS(appr_algo->maxM0_ == metadata.maxM0 && appr_algo->maxM_ == metadata.maxM, + "Layered HNSW M parameter mismatch"); + const auto allocation_elapsed_ms = elapsed_ms_since(allocation_start_time); + RAFT_LOG_INFO( + "Layered HNSW load: allocated hnswlib storage in %ld ms (base=%.2f GiB upper=%.2f GiB)", + allocation_elapsed_ms, + to_gib(metadata.n_rows * appr_algo->size_data_per_element_), + to_gib(metadata.upper_nodes_count * metadata.upper_link_row_bytes)); + + [[maybe_unused]] const auto num_threads = cuvs::core::omp::get_max_threads(); + + const size_t target_batch_bytes = 64 * 1024 * 1024; + const size_t dataset_row_bytes = metadata.dim * sizeof(T); + const size_t batch_size = std::max(1, target_batch_bytes / dataset_row_bytes); + auto dataset_buffer = + raft::make_host_matrix(static_cast(batch_size), metadata.dim); + + const auto base_start_time = std::chrono::steady_clock::now(); + std::chrono::steady_clock::duration dataset_read_time{}; + std::chrono::steady_clock::duration base_copy_time{}; + size_t dataset_bytes_read = 0; + for (size_t batch_start = 0; batch_start < metadata.n_rows; batch_start += batch_size) { + const auto current_batch_size = std::min(batch_size, metadata.n_rows - batch_start); + const auto current_dataset_bytes = current_batch_size * metadata.dim * sizeof(T); + auto batch_timer = std::chrono::steady_clock::now(); + cuvs::util::read_large_file(dataset_file.fd, + dataset_buffer.data_handle(), + current_dataset_bytes, + dataset_file.header_size + batch_start * metadata.dim * sizeof(T)); + dataset_read_time += std::chrono::steady_clock::now() - batch_timer; + dataset_bytes_read += current_dataset_bytes; + + bool link_list_allocation_failed = false; + batch_timer = std::chrono::steady_clock::now(); +#pragma omp parallel for num_threads(num_threads) reduction(|| : link_list_allocation_failed) + for (int64_t batch_idx = 0; batch_idx < static_cast(current_batch_size); ++batch_idx) { + const auto i = batch_start + static_cast(batch_idx); + auto level = static_cast(levels_u8[i]); + appr_algo->element_levels_[i] = level; + memcpy(appr_algo->getDataByInternalId(i), + dataset_buffer.data_handle() + batch_idx * metadata.dim, + appr_algo->data_size_); + *appr_algo->getExternalLabeLp(i) = static_cast(i); + if (level > 0) { + const auto link_list_size = appr_algo->size_links_per_element_ * level + 1; + appr_algo->linkLists_[i] = static_cast(malloc(link_list_size)); + if (appr_algo->linkLists_[i] == nullptr) { + link_list_allocation_failed = true; + continue; + } + } + } + if (link_list_allocation_failed) { + for (size_t i = 0; i < metadata.n_rows; ++i) { + free(appr_algo->linkLists_[i]); + appr_algo->linkLists_[i] = nullptr; + } + throw std::runtime_error("Not enough memory to allocate HNSW upper linklists"); + } + base_copy_time += std::chrono::steady_clock::now() - batch_timer; + log_deserialize_progress(current_dataset_bytes); + } + const auto base_elapsed_ms = elapsed_ms_since(base_start_time); + RAFT_LOG_INFO( + "Layered HNSW load: hnswlib data and levels initialized in %ld ms " + "(dataset read %.2f ms %.2f GiB %.2f GiB/s, copy %.2f ms)", + base_elapsed_ms, + elapsed_ms(dataset_read_time), + to_gib(dataset_bytes_read), + throughput_gib_per_s(dataset_bytes_read, dataset_read_time), + elapsed_ms(base_copy_time)); + + const auto base_topology_start_time = std::chrono::steady_clock::now(); + std::chrono::steady_clock::duration base_topology_read_time{}; + std::chrono::steady_clock::duration base_topology_copy_time{}; + size_t base_topology_bytes_read = 0; + const size_t base_topology_row_bytes = sizeof(uint32_t) + metadata.base_link_row_bytes; + const size_t base_topology_batch_size = + std::max(1, target_batch_bytes / base_topology_row_bytes); + std::vector base_node_buffer(base_topology_batch_size); + std::vector base_link_buffer(base_topology_batch_size * metadata.base_link_row_bytes); + for (size_t batch_start = 0; batch_start < metadata.n_rows; + batch_start += base_topology_batch_size) { + const auto current_batch_size = + std::min(base_topology_batch_size, metadata.n_rows - batch_start); + const auto current_base_topology_bytes = + current_batch_size * (sizeof(uint32_t) + metadata.base_link_row_bytes); + auto batch_timer = std::chrono::steady_clock::now(); + cuvs::util::read_large_file(artifact_fd, + base_node_buffer.data(), + current_batch_size * sizeof(uint32_t), + base_nodes_offset + batch_start * sizeof(uint32_t)); + cuvs::util::read_large_file(artifact_fd, + base_link_buffer.data(), + current_batch_size * metadata.base_link_row_bytes, + base_links_offset + batch_start * metadata.base_link_row_bytes); + base_topology_read_time += std::chrono::steady_clock::now() - batch_timer; + base_topology_bytes_read += current_base_topology_bytes; + + batch_timer = std::chrono::steady_clock::now(); + bool invalid_node_id = false; +#pragma omp parallel for num_threads(num_threads) reduction(|| : invalid_node_id) + for (int64_t batch_idx = 0; batch_idx < static_cast(current_batch_size); ++batch_idx) { + const auto node_id = static_cast(base_node_buffer[batch_idx]); + if (node_id >= metadata.n_rows) { + invalid_node_id = true; + continue; + } + auto ll0 = appr_algo->get_linklist0(node_id); + memcpy(ll0, + base_link_buffer.data() + batch_idx * metadata.base_link_row_bytes, + metadata.base_link_row_bytes); + } + RAFT_EXPECTS(!invalid_node_id, "Invalid base-layer node id in layered HNSW artifact"); + base_topology_copy_time += std::chrono::steady_clock::now() - batch_timer; + log_deserialize_progress(current_base_topology_bytes); + } + const auto base_topology_elapsed_ms = elapsed_ms_since(base_topology_start_time); + RAFT_LOG_INFO( + "Layered HNSW load: base-layer topology loaded in %ld ms " + "(read %.2f ms %.2f GiB %.2f GiB/s, scatter-copy %.2f ms)", + base_topology_elapsed_ms, + elapsed_ms(base_topology_read_time), + to_gib(base_topology_bytes_read), + throughput_gib_per_s(base_topology_bytes_read, base_topology_read_time), + elapsed_ms(base_topology_copy_time)); + + const auto upper_start_time = std::chrono::steady_clock::now(); + std::chrono::steady_clock::duration upper_read_time{}; + std::chrono::steady_clock::duration upper_copy_time{}; + size_t upper_bytes_read = 0; + for (const auto& layer : metadata.layers) { + const auto layer_row_bytes = sizeof(uint32_t) + metadata.upper_link_row_bytes; + const auto layer_batch_size = + std::min(std::max(1, target_batch_bytes / layer_row_bytes), layer.row_count); + std::vector node_buffer(layer_batch_size); + std::vector link_buffer(layer_batch_size * metadata.upper_link_row_bytes); + for (size_t batch_start = 0; batch_start < layer.row_count; batch_start += layer_batch_size) { + const auto current_batch_size = std::min(layer_batch_size, layer.row_count - batch_start); + const auto current_upper_bytes = + current_batch_size * (sizeof(uint32_t) + metadata.upper_link_row_bytes); + auto batch_timer = std::chrono::steady_clock::now(); + cuvs::util::read_large_file( + artifact_fd, + node_buffer.data(), + current_batch_size * sizeof(uint32_t), + upper_nodes_offset + (layer.node_offset + batch_start) * sizeof(uint32_t)); + cuvs::util::read_large_file( + artifact_fd, + link_buffer.data(), + current_batch_size * metadata.upper_link_row_bytes, + upper_links_offset + (layer.link_offset + batch_start) * metadata.upper_link_row_bytes); + upper_read_time += std::chrono::steady_clock::now() - batch_timer; + upper_bytes_read += current_upper_bytes; + batch_timer = std::chrono::steady_clock::now(); + bool invalid_node_id = false; + bool invalid_node_level = false; +#pragma omp parallel for num_threads(num_threads) \ + reduction(|| : invalid_node_id, invalid_node_level) + for (int64_t batch_idx = 0; batch_idx < static_cast(current_batch_size); + ++batch_idx) { + const auto node_id = static_cast(node_buffer[batch_idx]); + if (node_id >= metadata.n_rows) { + invalid_node_id = true; + continue; + } + if (layer.level > static_cast(levels_u8[node_id])) { + invalid_node_level = true; + continue; + } + auto ll = appr_algo->get_linklist(node_id, layer.level); + memcpy(ll, + link_buffer.data() + batch_idx * metadata.upper_link_row_bytes, + metadata.upper_link_row_bytes); + } + RAFT_EXPECTS(!invalid_node_id, "Invalid upper-layer node id in layered HNSW artifact"); + RAFT_EXPECTS(!invalid_node_level, + "Layered HNSW artifact references a node at an invalid upper level"); + upper_copy_time += std::chrono::steady_clock::now() - batch_timer; + log_deserialize_progress(current_upper_bytes); + } + } + const auto upper_elapsed_ms = elapsed_ms_since(upper_start_time); + RAFT_LOG_INFO( + "Layered HNSW load: upper layers loaded in %ld ms " + "(read %.2f ms %.2f GiB %.2f GiB/s, validate+copy %.2f ms)", + upper_elapsed_ms, + elapsed_ms(upper_read_time), + to_gib(upper_bytes_read), + throughput_gib_per_s(upper_bytes_read, upper_read_time), + elapsed_ms(upper_copy_time)); + + hnsw_index->set_index(std::move(appr_algo)); + RAFT_LOG_INFO("Layered HNSW load: total deserialize completed in %ld ms", + elapsed_ms_since(total_start_time)); + return hnsw_index; +} + template void deserialize(raft::resources const& res, const index_params& params, @@ -1647,6 +2892,11 @@ void deserialize(raft::resources const& res, cuvs::distance::DistanceType metric, index** idx) { + validate_output_format(params); + RAFT_EXPECTS(params.output_format == HnswOutputFormat::HNSWLIB, + "GRAPH_ONLY must be loaded with the two-filename hnsw::deserialize " + "overload"); + try { auto hnsw_index = std::make_unique>(dim, metric, params.hierarchy); auto appr_algo = std::make_unique::type>>( @@ -1663,6 +2913,16 @@ void deserialize(raft::resources const& res, } } +template +void deserialize(raft::resources const& res, + const std::string& graph_filename, + const std::string& dataset_filename, + index** idx) +{ + auto hnsw_index = deserialize_layered_hnswlib(res, graph_filename, dataset_filename); + *idx = hnsw_index.release(); +} + /** * @brief Build an HNSW index on the GPU using CAGRA graph building algorithm * @@ -1678,6 +2938,19 @@ std::unique_ptr> build(raft::resources const& res, raft::host_matrix_view dataset) { common::nvtx::range fun_scope("hnsw::build"); + validate_output_format(params); + + // GRAPH_ONLY materializes the layered index from disk-backed ACE artifacts, so it + // requires a GPU hierarchy and ACE disk mode with a build directory. Validate up front; this + // also forces the CAGRA build below onto the ACE disk path (use_ace becomes true because ACE + // params are set). + if (params.output_format == HnswOutputFormat::GRAPH_ONLY) { + RAFT_EXPECTS(std::holds_alternative(params.graph_build_params), + "GRAPH_ONLY requires ACE parameters to be configured"); + const auto& ace = std::get(params.graph_build_params); + RAFT_EXPECTS(ace.use_disk, "GRAPH_ONLY requires ACE disk mode (ace_params.use_disk = true)"); + RAFT_EXPECTS(!ace.build_dir.empty(), "GRAPH_ONLY requires ace_params.build_dir to be set"); + } cuvs::neighbors::cagra::index_params cagra_params = cagra::index_params::from_hnsw_params(dataset.extents(), @@ -1736,11 +3009,7 @@ std::unique_ptr> build(raft::resources const& res, RAFT_LOG_INFO("hnsw::build - Converting CAGRA index to HNSW format"); - return from_cagra( - res, - params, - ace_host_index, - ace_host_index.dataset_fd().has_value() ? std::nullopt : std::make_optional(dataset)); + return from_cagra(res, params, ace_host_index, std::make_optional(dataset)); } } // namespace cuvs::neighbors::hnsw::detail diff --git a/cpp/src/neighbors/detail/hnsw_layered_format.hpp b/cpp/src/neighbors/detail/hnsw_layered_format.hpp new file mode 100644 index 0000000000..cce5b12b47 --- /dev/null +++ b/cpp/src/neighbors/detail/hnsw_layered_format.hpp @@ -0,0 +1,74 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +// On-disk layout of the layered HNSW artifact. + +#include +#include + +namespace cuvs::neighbors::hnsw::detail { + +enum class layered_hnsw_dtype : uint32_t { + unknown = 0, + float32 = 1, + float16 = 2, + uint8 = 3, + int8 = 4, +}; + +// The layered HNSW artifact begins with this fixed-size POD header, immediately followed by +// `num_layers` layered_hnsw_layer_descriptor records and then the payload sections. +struct layered_hnsw_file_header { + char magic[32]; + uint64_t n_rows; + uint64_t dim; + uint64_t M; + uint64_t maxM; + uint64_t maxM0; + uint64_t ef_construction; + uint64_t base_degree; + uint64_t levels_bytes; + uint64_t base_nodes_bytes; + uint64_t base_link_row_bytes; + uint64_t base_links_bytes; + uint64_t upper_nodes_count; + uint64_t upper_nodes_bytes; + uint64_t upper_link_row_bytes; + uint64_t upper_links_bytes; + double mult; + uint32_t version; + // Element type used to construct the graph. Attached vectors may use a different type. + uint32_t construction_dtype; + uint32_t metric; + uint32_t num_layers; + int32_t maxlevel; + int32_t enterpoint_node; + uint32_t reserved0; + uint32_t reserved1; +}; +static_assert(sizeof(layered_hnsw_file_header) == 192, + "layered_hnsw_file_header must keep a fixed 192-byte on-disk layout"); +inline constexpr size_t layered_hnsw_file_header_version_offset = + offsetof(layered_hnsw_file_header, version); +inline constexpr size_t layered_hnsw_file_header_construction_dtype_offset = + offsetof(layered_hnsw_file_header, construction_dtype); + +struct layered_hnsw_layer_descriptor { + uint64_t level; + uint64_t row_count; + uint64_t degree; + uint64_t node_offset; + uint64_t link_offset; +}; +static_assert(sizeof(layered_hnsw_layer_descriptor) == 40, + "layered_hnsw_layer_descriptor must keep a fixed 40-byte on-disk layout"); + +constexpr const char* layered_hnsw_magic = "CUVS_HNSW_LAYERED"; +constexpr uint32_t layered_hnsw_version = 1; +constexpr size_t layered_hnsw_alignment = 64; + +} // namespace cuvs::neighbors::hnsw::detail diff --git a/cpp/src/neighbors/hnsw.cpp b/cpp/src/neighbors/hnsw.cpp index 3958366311..026f8a9bdd 100644 --- a/cpp/src/neighbors/hnsw.cpp +++ b/cpp/src/neighbors/hnsw.cpp @@ -157,6 +157,13 @@ CUVS_INST_HNSW_SEARCH(int8_t); index** idx) \ { \ detail::deserialize(res, params, filename, dim, metric, idx); \ + } \ + void deserialize(raft::resources const& res, \ + const std::string& graph_filename, \ + const std::string& dataset_filename, \ + index** idx) \ + { \ + detail::deserialize(res, graph_filename, dataset_filename, idx); \ } CUVS_INST_HNSW_SERIALIZE(float); diff --git a/cpp/src/neighbors/ivf_flat/ivf_flat_serialize.cuh b/cpp/src/neighbors/ivf_flat/ivf_flat_serialize.cuh index 9060f477de..0dfcf5d552 100644 --- a/cpp/src/neighbors/ivf_flat/ivf_flat_serialize.cuh +++ b/cpp/src/neighbors/ivf_flat/ivf_flat_serialize.cuh @@ -45,7 +45,7 @@ void serialize(raft::resources const& handle, Output& os, const index& RAFT_LOG_DEBUG( "Saving IVF-Flat index, size %zu, dim %u", static_cast(index_.size()), index_.dim()); - std::string dtype_string = raft::numpy_serializer::get_numpy_dtype().to_string(); + std::string dtype_string = cuvs::util::detail::numpy_dtype_string(); dtype_string.resize(4); os << dtype_string; diff --git a/cpp/src/neighbors/ivf_sq/ivf_sq_serialize.cuh b/cpp/src/neighbors/ivf_sq/ivf_sq_serialize.cuh index 499e678c03..72ff28cf21 100644 --- a/cpp/src/neighbors/ivf_sq/ivf_sq_serialize.cuh +++ b/cpp/src/neighbors/ivf_sq/ivf_sq_serialize.cuh @@ -12,7 +12,6 @@ #include #include -#include #include #include #include @@ -29,7 +28,7 @@ void serialize(raft::resources const& handle, Output& os, const index& in RAFT_LOG_DEBUG( "Saving IVF-SQ index, size %zu, dim %u", static_cast(index_.size()), index_.dim()); - std::string dtype_string = raft::detail::numpy_serializer::get_numpy_dtype().to_string(); + std::string dtype_string = cuvs::util::detail::numpy_dtype_string(); dtype_string.resize(4); os << dtype_string; diff --git a/cpp/src/neighbors/mg/snmg.cuh b/cpp/src/neighbors/mg/snmg.cuh index 9e66f58898..d93dcb5163 100644 --- a/cpp/src/neighbors/mg/snmg.cuh +++ b/cpp/src/neighbors/mg/snmg.cuh @@ -25,6 +25,7 @@ #include #include #include +#include #include @@ -785,7 +786,7 @@ void serialize(const raft::resources& clique, { cuvs::util::kvikio_ofstream of(filename); - std::string dtype_string = raft::numpy_serializer::get_numpy_dtype().to_string(); + std::string dtype_string = cuvs::util::detail::numpy_dtype_string(); dtype_string.resize(4); of << dtype_string; diff --git a/cpp/src/util/serialize_validation.hpp b/cpp/src/util/serialize_validation.hpp index b30dd0d390..0e8faea7c6 100644 --- a/cpp/src/util/serialize_validation.hpp +++ b/cpp/src/util/serialize_validation.hpp @@ -1,13 +1,13 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once #include #include +#include -#include #include #include @@ -44,7 +44,7 @@ inline bool validate_serialized_dtype(const char* dtype_prefix, std::size_t dtyp { if (dtype_prefix == nullptr || dtype_prefix_size != 4) { return false; } - auto expected_dtype = raft::detail::numpy_serializer::get_numpy_dtype().to_string(); + auto expected_dtype = cuvs::util::detail::numpy_dtype_string(); expected_dtype.resize(dtype_prefix_size, '\0'); return std::equal(dtype_prefix, dtype_prefix + dtype_prefix_size, expected_dtype.begin()); diff --git a/cpp/tests/neighbors/ann_hnsw_ace.cuh b/cpp/tests/neighbors/ann_hnsw_ace.cuh index f953da0c55..90a3e46caf 100644 --- a/cpp/tests/neighbors/ann_hnsw_ace.cuh +++ b/cpp/tests/neighbors/ann_hnsw_ace.cuh @@ -9,18 +9,25 @@ #include #include +#include "../../src/neighbors/detail/hnsw_layered_format.hpp" + #include #include #include +#include #include +#include #include #include #include +#include #include #include #include +#include #include +#include #include @@ -437,6 +444,298 @@ class AnnHnswAceTest : public ::testing::TestWithParam { std::filesystem::remove_all(temp_dir); } + void testHnswAceLayeredBuildDeserializeSearch() + { + size_t queries_size = ps.n_queries * ps.k; + std::vector indexes_naive(queries_size); + std::vector distances_naive(queries_size); + + { + rmm::device_uvector distances_naive_dev(queries_size, stream_); + rmm::device_uvector indexes_naive_dev(queries_size, stream_); + + cuvs::neighbors::naive_knn(handle_, + distances_naive_dev.data(), + indexes_naive_dev.data(), + search_queries.data(), + database_dev.data(), + ps.n_queries, + ps.n_rows, + ps.dim, + ps.k, + ps.metric); + raft::update_host(distances_naive.data(), distances_naive_dev.data(), queries_size, stream_); + raft::update_host(indexes_naive.data(), indexes_naive_dev.data(), queries_size, stream_); + raft::resource::sync_stream(handle_); + } + + std::string temp_dir = std::string("/tmp/cuvs_hnsw_ace_layered_test_") + + std::to_string(std::time(nullptr)) + "_" + + std::to_string(reinterpret_cast(this)); + std::filesystem::create_directories(temp_dir); + struct temp_dir_cleanup { + std::string path; + ~temp_dir_cleanup() + { + std::error_code ec; + std::filesystem::remove_all(path, ec); + } + } cleanup{temp_dir}; + + auto database_host = raft::make_host_matrix(ps.n_rows, ps.dim); + raft::copy(database_host.data_handle(), database_dev.data(), ps.n_rows * ps.dim, stream_); + auto queries_host = raft::make_host_matrix(ps.n_queries, ps.dim); + raft::copy(queries_host.data_handle(), search_queries.data(), ps.n_queries * ps.dim, stream_); + raft::resource::sync_stream(handle_); + + const auto dataset_file = (std::filesystem::path(temp_dir) / "dataset.npy").string(); + auto [dataset_fd, dataset_header_size] = cuvs::util::create_numpy_file( + dataset_file, {static_cast(ps.n_rows), static_cast(ps.dim)}); + cuvs::util::write_large_file(dataset_fd, + database_host.data_handle(), + static_cast(ps.n_rows) * ps.dim * sizeof(DataT), + dataset_header_size); + + hnsw::index_params hnsw_params; + hnsw_params.metric = ps.metric; + hnsw_params.hierarchy = hnsw::HnswHierarchy::GPU; + hnsw_params.output_format = hnsw::HnswOutputFormat::GRAPH_ONLY; + hnsw_params.M = 32; + hnsw_params.ef_construction = ps.ef_construction; + + auto ace_params = graph_build_params::ace_params(); + ace_params.npartitions = ps.npartitions; + ace_params.build_dir = temp_dir; + ace_params.use_disk = true; + ace_params.max_host_memory_gb = ps.max_host_memory_gb; + ace_params.max_gpu_memory_gb = ps.max_gpu_memory_gb; + hnsw_params.graph_build_params = ace_params; + + auto invalid_params = hnsw_params; + invalid_params.hierarchy = hnsw::HnswHierarchy::CPU; + EXPECT_THROW( + hnsw::build(handle_, invalid_params, raft::make_const_mdspan(database_host.view())), + std::exception); + + auto hnsw_index = + hnsw::build(handle_, hnsw_params, raft::make_const_mdspan(database_host.view())); + ASSERT_NE(hnsw_index, nullptr); + EXPECT_EQ(hnsw_index->hierarchy(), hnsw::HnswHierarchy::GPU); + EXPECT_EQ(hnsw_index->output_format(), hnsw::HnswOutputFormat::GRAPH_ONLY); + + const auto artifact_path = hnsw_index->file_path(); + ASSERT_FALSE(artifact_path.empty()); + ASSERT_TRUE(std::filesystem::is_regular_file(artifact_path)); + EXPECT_EQ(std::filesystem::path(artifact_path).filename().string(), "hnsw_index.cuvs"); + size_t cuvs_artifact_count = 0; + for (const auto& entry : std::filesystem::directory_iterator(temp_dir)) { + if (entry.path().extension() == ".cuvs") { ++cuvs_artifact_count; } + } + EXPECT_EQ(cuvs_artifact_count, 1); + EXPECT_FALSE(std::filesystem::exists(std::filesystem::path(temp_dir) / "layered_hnsw")); + + auto indexes_hnsw_host = raft::make_host_matrix(ps.n_queries, ps.k); + auto distances_hnsw_host = raft::make_host_matrix(ps.n_queries, ps.k); + + hnsw::search_params search_params; + search_params.ef = std::max(ps.ef_construction, ps.k * 2); + search_params.num_threads = 1; + + EXPECT_THROW(hnsw::search(handle_, + search_params, + *hnsw_index, + queries_host.view(), + indexes_hnsw_host.view(), + distances_hnsw_host.view()), + std::exception); + + hnsw::index* deserialized_index = nullptr; + hnsw::deserialize(handle_, artifact_path, dataset_file, &deserialized_index); + ASSERT_NE(deserialized_index, nullptr); + std::unique_ptr> deserialized_guard(deserialized_index); + EXPECT_EQ(deserialized_guard->hierarchy(), hnsw::HnswHierarchy::GPU); + EXPECT_EQ(deserialized_guard->output_format(), hnsw::HnswOutputFormat::GRAPH_ONLY); + + hnsw::search(handle_, + search_params, + *deserialized_guard, + queries_host.view(), + indexes_hnsw_host.view(), + distances_hnsw_host.view()); + + std::vector indexes_hnsw_converted(queries_size); + std::vector distances_hnsw(queries_size); + for (size_t i = 0; i < queries_size; i++) { + indexes_hnsw_converted[i] = static_cast(indexes_hnsw_host.data_handle()[i]); + distances_hnsw[i] = distances_hnsw_host.data_handle()[i]; + } + + EXPECT_TRUE(cuvs::neighbors::eval_neighbours(indexes_naive, + indexes_hnsw_converted, + distances_naive, + distances_hnsw, + ps.n_queries, + ps.k, + 0.003, + ps.min_recall)) + << "Layered HNSW deserialize and search failed recall check"; + + if constexpr (std::is_same_v) { + auto database_float = raft::make_host_matrix(ps.n_rows, ps.dim); + auto queries_float = raft::make_host_matrix(ps.n_queries, ps.dim); + std::transform(database_host.data_handle(), + database_host.data_handle() + ps.n_rows * ps.dim, + database_float.data_handle(), + [](int8_t value) { return static_cast(value); }); + std::transform(queries_host.data_handle(), + queries_host.data_handle() + ps.n_queries * ps.dim, + queries_float.data_handle(), + [](int8_t value) { return static_cast(value); }); + + const auto float_dataset_file = + (std::filesystem::path(temp_dir) / "dataset_float.npy").string(); + auto [float_dataset_fd, float_dataset_header_size] = cuvs::util::create_numpy_file( + float_dataset_file, {static_cast(ps.n_rows), static_cast(ps.dim)}); + cuvs::util::write_large_file(float_dataset_fd, + database_float.data_handle(), + static_cast(ps.n_rows) * ps.dim * sizeof(float), + float_dataset_header_size); + + hnsw::index* mixed_precision_index = nullptr; + hnsw::deserialize(handle_, artifact_path, float_dataset_file, &mixed_precision_index); + ASSERT_NE(mixed_precision_index, nullptr); + std::unique_ptr> mixed_precision_guard(mixed_precision_index); + + hnsw::search(handle_, + search_params, + *mixed_precision_guard, + queries_float.view(), + indexes_hnsw_host.view(), + distances_hnsw_host.view()); + + for (size_t i = 0; i < queries_size; i++) { + indexes_hnsw_converted[i] = static_cast(indexes_hnsw_host.data_handle()[i]); + distances_hnsw[i] = distances_hnsw_host.data_handle()[i]; + } + EXPECT_TRUE(cuvs::neighbors::eval_neighbours(indexes_naive, + indexes_hnsw_converted, + distances_naive, + distances_hnsw, + ps.n_queries, + ps.k, + 0.003, + ps.min_recall)) + << "Layered HNSW int8 topology with float vectors failed recall check"; + } + + const auto copied_artifact = + (std::filesystem::path(temp_dir) / "copied_layered" / "hnsw_index.cuvs").string(); + hnsw::serialize(handle_, copied_artifact, *hnsw_index); + EXPECT_TRUE(std::filesystem::is_regular_file(copied_artifact)); + size_t copied_file_count = 0; + for (const auto& entry : std::filesystem::directory_iterator( + std::filesystem::path(copied_artifact).parent_path())) { + if (entry.is_regular_file()) { ++copied_file_count; } + } + EXPECT_EQ(copied_file_count, 1); + + hnsw::index* copied_index = nullptr; + hnsw::deserialize(handle_, copied_artifact, dataset_file, &copied_index); + ASSERT_NE(copied_index, nullptr); + std::unique_ptr> copied_guard(copied_index); + + hnsw::search(handle_, + search_params, + *copied_guard, + queries_host.view(), + indexes_hnsw_host.view(), + distances_hnsw_host.view()); + + for (size_t i = 0; i < queries_size; i++) { + indexes_hnsw_converted[i] = static_cast(indexes_hnsw_host.data_handle()[i]); + distances_hnsw[i] = distances_hnsw_host.data_handle()[i]; + } + EXPECT_TRUE(cuvs::neighbors::eval_neighbours(indexes_naive, + indexes_hnsw_converted, + distances_naive, + distances_hnsw, + ps.n_queries, + ps.k, + 0.003, + ps.min_recall)) + << "Copied layered HNSW artifact deserialize and search failed recall check"; + + const auto bad_artifact = + (std::filesystem::path(temp_dir) / "bad_layered" / "bad_magic.cuvs").string(); + std::filesystem::create_directories(std::filesystem::path(bad_artifact).parent_path()); + { + std::ofstream bad_file(bad_artifact, std::ios::binary); + std::array bad_bytes{}; + bad_file.write(bad_bytes.data(), bad_bytes.size()); + } + hnsw::index* bad_index = nullptr; + EXPECT_THROW(hnsw::deserialize(handle_, bad_artifact, dataset_file, &bad_index), + std::exception); + + const auto bad_version_artifact = + (std::filesystem::path(temp_dir) / "bad_layered" / "bad_version.cuvs").string(); + std::filesystem::copy_file( + copied_artifact, bad_version_artifact, std::filesystem::copy_options::overwrite_existing); + { + std::fstream bad_version_file(bad_version_artifact, + std::ios::in | std::ios::out | std::ios::binary); + const uint32_t bad_version = 999; + bad_version_file.seekp(static_cast( + cuvs::neighbors::hnsw::detail::layered_hnsw_file_header_version_offset)); + bad_version_file.write(reinterpret_cast(&bad_version), sizeof(bad_version)); + } + hnsw::index* bad_version_index = nullptr; + EXPECT_THROW(hnsw::deserialize(handle_, bad_version_artifact, dataset_file, &bad_version_index), + std::exception); + + const auto bad_dtype_artifact = + (std::filesystem::path(temp_dir) / "bad_layered" / "bad_construction_dtype.cuvs").string(); + std::filesystem::copy_file( + copied_artifact, bad_dtype_artifact, std::filesystem::copy_options::overwrite_existing); + { + std::fstream bad_dtype_file(bad_dtype_artifact, + std::ios::in | std::ios::out | std::ios::binary); + const uint32_t bad_dtype = 999; + bad_dtype_file.seekp(static_cast( + cuvs::neighbors::hnsw::detail::layered_hnsw_file_header_construction_dtype_offset)); + bad_dtype_file.write(reinterpret_cast(&bad_dtype), sizeof(bad_dtype)); + } + hnsw::index* bad_dtype_index = nullptr; + EXPECT_THROW(hnsw::deserialize(handle_, bad_dtype_artifact, dataset_file, &bad_dtype_index), + std::exception); + + hnsw::index* missing_dataset_index = nullptr; + EXPECT_THROW(hnsw::deserialize(handle_, copied_artifact, "", &missing_dataset_index), + std::exception); + + const auto truncated_artifact = + (std::filesystem::path(temp_dir) / "bad_layered" / "truncated.cuvs").string(); + std::filesystem::copy_file( + copied_artifact, truncated_artifact, std::filesystem::copy_options::overwrite_existing); + std::filesystem::resize_file(truncated_artifact, 128); + hnsw::index* truncated_index = nullptr; + EXPECT_THROW(hnsw::deserialize(handle_, truncated_artifact, dataset_file, &truncated_index), + std::exception); + + const auto wrong_dataset_file = + (std::filesystem::path(temp_dir) / "bad_layered" / "wrong_dataset.npy").string(); + auto [wrong_dataset_fd, wrong_dataset_header_size] = cuvs::util::create_numpy_file( + wrong_dataset_file, {static_cast(ps.n_rows - 1), static_cast(ps.dim)}); + cuvs::util::write_large_file(wrong_dataset_fd, + database_host.data_handle(), + static_cast(ps.n_rows - 1) * ps.dim * sizeof(DataT), + wrong_dataset_header_size); + hnsw::index* wrong_dataset_index = nullptr; + EXPECT_THROW( + hnsw::deserialize(handle_, copied_artifact, wrong_dataset_file, &wrong_dataset_index), + std::exception); + } + void testHnswAceRejectsTooManyPartitions() { auto database_host = raft::make_host_matrix(ps.n_rows, ps.dim); @@ -704,6 +1003,23 @@ inline std::vector generate_hnsw_ace_memory_fallback_inputs() }; } +inline std::vector generate_hnsw_ace_layered_inputs() +{ + return { + {10, // n_queries + 5000, // n_rows + 64, // dim + 10, // k + 2, // npartitions + 100, // ef_construction + true, // use_disk + cuvs::distance::DistanceType::L2Expanded, + 0.9, // min_recall + 0.0, // max_host_memory_gb + 0.0} // max_gpu_memory_gb + }; +} + inline std::vector generate_hnsw_ace_invalid_partition_inputs() { return {{10, @@ -739,6 +1055,7 @@ inline std::vector generate_hnsw_inmem_spill_inputs() const std::vector hnsw_ace_inputs = generate_hnsw_ace_inputs(); const std::vector hnsw_ace_memory_fallback_inputs = generate_hnsw_ace_memory_fallback_inputs(); +const std::vector hnsw_ace_layered_inputs = generate_hnsw_ace_layered_inputs(); const std::vector hnsw_ace_invalid_partition_inputs = generate_hnsw_ace_invalid_partition_inputs(); const std::vector hnsw_inmem_spill_inputs = generate_hnsw_inmem_spill_inputs(); diff --git a/cpp/tests/neighbors/ann_hnsw_ace/test_float_uint32_t.cu b/cpp/tests/neighbors/ann_hnsw_ace/test_float_uint32_t.cu index 4efa96b2a5..6cadb07646 100644 --- a/cpp/tests/neighbors/ann_hnsw_ace/test_float_uint32_t.cu +++ b/cpp/tests/neighbors/ann_hnsw_ace/test_float_uint32_t.cu @@ -53,6 +53,16 @@ INSTANTIATE_TEST_CASE_P(AnnHnswAceMemoryFallbackTest, AnnHnswAceMemoryFallbackTest_float, ::testing::ValuesIn(hnsw_ace_memory_fallback_inputs)); +typedef AnnHnswAceTest AnnHnswAceLayeredTest_float; +TEST_P(AnnHnswAceLayeredTest_float, AnnHnswAceLayeredBuildDeserializeSearch) +{ + this->testHnswAceLayeredBuildDeserializeSearch(); +} + +INSTANTIATE_TEST_CASE_P(AnnHnswAceLayeredTest, + AnnHnswAceLayeredTest_float, + ::testing::ValuesIn(hnsw_ace_layered_inputs)); + // Test for in-memory CAGRA -> HNSW disk-spill conversion typedef AnnHnswAceTest AnnHnswInmemSpillTest_float; TEST_P(AnnHnswInmemSpillTest_float, AnnHnswFromCagraInmemSpill) diff --git a/cpp/tests/neighbors/ann_hnsw_ace/test_half_uint32_t.cu b/cpp/tests/neighbors/ann_hnsw_ace/test_half_uint32_t.cu index 29cba2fef3..67d72c81d8 100644 --- a/cpp/tests/neighbors/ann_hnsw_ace/test_half_uint32_t.cu +++ b/cpp/tests/neighbors/ann_hnsw_ace/test_half_uint32_t.cu @@ -43,6 +43,16 @@ INSTANTIATE_TEST_CASE_P(AnnHnswAceMemoryFallbackTest, AnnHnswAceMemoryFallbackTest_half, ::testing::ValuesIn(hnsw_ace_memory_fallback_inputs)); +typedef AnnHnswAceTest AnnHnswAceLayeredTest_half; +TEST_P(AnnHnswAceLayeredTest_half, AnnHnswAceLayeredBuildDeserializeSearch) +{ + this->testHnswAceLayeredBuildDeserializeSearch(); +} + +INSTANTIATE_TEST_CASE_P(AnnHnswAceLayeredTest, + AnnHnswAceLayeredTest_half, + ::testing::ValuesIn(hnsw_ace_layered_inputs)); + // Test for in-memory CAGRA -> HNSW disk-spill conversion typedef AnnHnswAceTest AnnHnswInmemSpillTest_half; TEST_P(AnnHnswInmemSpillTest_half, AnnHnswFromCagraInmemSpill) diff --git a/cpp/tests/neighbors/ann_hnsw_ace/test_int8_t_uint32_t.cu b/cpp/tests/neighbors/ann_hnsw_ace/test_int8_t_uint32_t.cu index 5c203096c3..52202d7a22 100644 --- a/cpp/tests/neighbors/ann_hnsw_ace/test_int8_t_uint32_t.cu +++ b/cpp/tests/neighbors/ann_hnsw_ace/test_int8_t_uint32_t.cu @@ -45,6 +45,16 @@ INSTANTIATE_TEST_CASE_P(AnnHnswAceMemoryFallbackTest, AnnHnswAceMemoryFallbackTest_int8_t, ::testing::ValuesIn(hnsw_ace_memory_fallback_inputs)); +typedef AnnHnswAceTest AnnHnswAceLayeredTest_int8_t; +TEST_P(AnnHnswAceLayeredTest_int8_t, AnnHnswAceLayeredBuildDeserializeSearch) +{ + this->testHnswAceLayeredBuildDeserializeSearch(); +} + +INSTANTIATE_TEST_CASE_P(AnnHnswAceLayeredTest, + AnnHnswAceLayeredTest_int8_t, + ::testing::ValuesIn(hnsw_ace_layered_inputs)); + // Test for in-memory CAGRA -> HNSW disk-spill conversion typedef AnnHnswAceTest AnnHnswInmemSpillTest_int8_t; TEST_P(AnnHnswInmemSpillTest_int8_t, AnnHnswFromCagraInmemSpill) diff --git a/cpp/tests/neighbors/ann_hnsw_ace/test_uint8_t_uint32_t.cu b/cpp/tests/neighbors/ann_hnsw_ace/test_uint8_t_uint32_t.cu index f9c754e7d6..a0825c0c9d 100644 --- a/cpp/tests/neighbors/ann_hnsw_ace/test_uint8_t_uint32_t.cu +++ b/cpp/tests/neighbors/ann_hnsw_ace/test_uint8_t_uint32_t.cu @@ -45,6 +45,16 @@ INSTANTIATE_TEST_CASE_P(AnnHnswAceMemoryFallbackTest, AnnHnswAceMemoryFallbackTest_uint8_t, ::testing::ValuesIn(hnsw_ace_memory_fallback_inputs)); +typedef AnnHnswAceTest AnnHnswAceLayeredTest_uint8_t; +TEST_P(AnnHnswAceLayeredTest_uint8_t, AnnHnswAceLayeredBuildDeserializeSearch) +{ + this->testHnswAceLayeredBuildDeserializeSearch(); +} + +INSTANTIATE_TEST_CASE_P(AnnHnswAceLayeredTest, + AnnHnswAceLayeredTest_uint8_t, + ::testing::ValuesIn(hnsw_ace_layered_inputs)); + // Test for in-memory CAGRA -> HNSW disk-spill conversion typedef AnnHnswAceTest AnnHnswInmemSpillTest_uint8_t; TEST_P(AnnHnswInmemSpillTest_uint8_t, AnnHnswFromCagraInmemSpill) diff --git a/examples/cpp/CMakeLists.txt b/examples/cpp/CMakeLists.txt index f4dffad3a9..657b30d92e 100644 --- a/examples/cpp/CMakeLists.txt +++ b/examples/cpp/CMakeLists.txt @@ -46,6 +46,7 @@ add_executable(CAGRA_HNSW_ACE_EXAMPLE src/cagra_hnsw_ace_example.cu) add_executable(CAGRA_PERSISTENT_EXAMPLE src/cagra_persistent_example.cu) add_executable(DYNAMIC_BATCHING_EXAMPLE src/dynamic_batching_example.cu) add_executable(HNSW_ACE_EXAMPLE src/hnsw_ace_example.cu) +add_executable(HNSW_ACE_LAYERED_EXAMPLE src/hnsw_ace_layered_example.cu) add_executable(HNSW_OPENAI_EXAMPLE src/hnsw_openai_example.cu) add_executable(IVF_FLAT_EXAMPLE src/ivf_flat_example.cu) add_executable(IVF_PQ_EXAMPLE src/ivf_pq_example.cu) @@ -74,8 +75,10 @@ target_link_libraries( DYNAMIC_BATCHING_EXAMPLE PRIVATE cuvs::cuvs $ Threads::Threads ) target_link_libraries(HNSW_ACE_EXAMPLE PRIVATE cuvs::cuvs $) +target_link_libraries( + HNSW_ACE_LAYERED_EXAMPLE PRIVATE cuvs::cuvs $ +) target_link_libraries(HNSW_OPENAI_EXAMPLE PRIVATE cuvs::cuvs $) - target_link_libraries(IVF_PQ_EXAMPLE PRIVATE cuvs::cuvs $) target_link_libraries(IVF_FLAT_EXAMPLE PRIVATE cuvs::cuvs $) target_link_libraries(VAMANA_EXAMPLE PRIVATE cuvs::cuvs $) diff --git a/examples/cpp/src/hnsw_ace_layered_example.cu b/examples/cpp/src/hnsw_ace_layered_example.cu new file mode 100644 index 0000000000..14de426605 --- /dev/null +++ b/examples/cpp/src/hnsw_ace_layered_example.cu @@ -0,0 +1,261 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +// Large graphs exceeding the memory capacity can be built using the Augmented Core Extraction (ACE) +// algorithm, which partitions the dataset. The resulting HNSW index is too large to fit in memory +// as well. Thus, the index needs to be transferred to a search server with enough memory. +// +// HnswOutputFormat::GRAPH_ONLY builds a GPU hierarchy as a graph-only HNSW artifact on disk. +// It emits one graph artifact, hnsw_index.cuvs. The dataset remains separate and does not +// need to be transferred to the search server, which typically has the dataset locally. +// +// This example demonstrates how to build a graph-only HNSW artifact with ACE: +// +// 1. Optionally quantize the dataset to int8 for graph construction. +// 2. Build a single-file graph-only HNSW artifact with ACE using hnsw::build. +// 3. Attach the original float dataset using the two-filename hnsw::deserialize overload. +// 4. Search the in-memory float HNSW index with the original float queries. +// +// Layered-on-disk layout: +// +// index_dir/hnsw_index.cuvs +// fixed header + layer descriptors +// levels: uint8 [N], max HNSW level for each original row id +// base nodes + base links: uint32 node ids with hnswlib-ready link rows +// upper nodes + upper links: hnswlib-ready upper-layer topology +// +// The transferred index artifact is graph-only. The dataset filename is passed separately to +// deserialize. The loader supports row-major .npy files and type-specific ANN benchmark binary +// files (.fbin, .f16bin/.fp16.fbin, .u8bin, and .i8bin). This example writes a local .npy dataset +// only to make the demo self-contained. +// +// Layer 0 node IDs and neighbor IDs are original dataset row IDs. Upper layers are generated with +// the same level/order/KNN logic as serialize_to_hnswlib_from_disk, then stored as hnswlib-ready +// link rows so deserialization does no graph remapping or link padding on the search node. + +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +#include +#include +#include + +#include + +#include "common.cuh" + +// When 1, scalar-quantize the float dataset to int8. +#define HNSW_ACE_LAYERED_USE_QUANTIZATION 1 + +namespace { + +constexpr const char* kBuildDir = "/tmp/hnsw_ace_layered"; + +template +std::string write_local_dataset(raft::host_matrix_view dataset, + const std::string& path) +{ + auto [fd, header_size] = cuvs::util::create_numpy_file( + path, {static_cast(dataset.extent(0)), static_cast(dataset.extent(1))}); + cuvs::util::write_large_file( + fd, dataset.data_handle(), dataset.extent(0) * dataset.extent(1) * sizeof(T), header_size); + return path; +} + +auto quantize_dataset(raft::device_resources const& dev_resources, + raft::host_matrix_view dataset_float) + -> raft::host_matrix +{ + std::cout << " quantize_dataset: training scalar quantizer (float -> int8)" << std::endl; + cuvs::preprocessing::quantize::scalar::params qp; + auto quantizer = cuvs::preprocessing::quantize::scalar::train(dev_resources, qp, dataset_float); + + auto dataset_i8 = + raft::make_host_matrix(dataset_float.extent(0), dataset_float.extent(1)); + cuvs::preprocessing::quantize::scalar::transform( + dev_resources, quantizer, dataset_float, dataset_i8.view()); + return dataset_i8; +} + +auto make_hnsw_ace_params(const std::string& build_dir) -> cuvs::neighbors::hnsw::index_params +{ + using namespace cuvs::neighbors; + + hnsw::index_params hnsw_params; + hnsw_params.metric = cuvs::distance::DistanceType::L2Expanded; + hnsw_params.hierarchy = hnsw::HnswHierarchy::GPU; + hnsw_params.output_format = hnsw::HnswOutputFormat::GRAPH_ONLY; + hnsw_params.M = 32; + hnsw_params.ef_construction = 120; + + auto ace_params = hnsw::graph_build_params::ace_params(); + ace_params.npartitions = 4; + ace_params.use_disk = true; + ace_params.build_dir = build_dir; + hnsw_params.graph_build_params = ace_params; + + return hnsw_params; +} + +template +auto hnsw_build(raft::device_resources const& dev_resources, + const cuvs::neighbors::hnsw::index_params& hnsw_params, + raft::host_matrix_view dataset) -> std::string +{ + using namespace cuvs::neighbors; + + auto hnsw_index = hnsw::build(dev_resources, hnsw_params, dataset); + const auto artifact_path = hnsw_index->file_path(); + if (artifact_path.empty()) { + throw std::runtime_error("Expected layered HNSW build to return an artifact path."); + } + std::cout << " hnsw_build: layered artifact written to " << artifact_path << std::endl; + return artifact_path; +} + +template +auto hnsw_deserialize(raft::device_resources const& dev_resources, + const std::string& artifact_path, + const std::string& dataset_path) + -> std::unique_ptr> +{ + using namespace cuvs::neighbors; + + hnsw::index* deserialized_index = nullptr; + hnsw::deserialize(dev_resources, artifact_path, dataset_path, &deserialized_index); + return std::unique_ptr>(deserialized_index); +} + +template +void hnsw_search(raft::device_resources const& dev_resources, + const cuvs::neighbors::hnsw::index& hnsw_index, + raft::host_matrix_view queries, + int64_t topk = 12) +{ + using namespace cuvs::neighbors; + + const int64_t n_queries = queries.extent(0); + auto indices_hnsw_host = raft::make_host_matrix(n_queries, topk); + auto distances_hnsw_host = raft::make_host_matrix(n_queries, topk); + + hnsw::search_params search_params; + search_params.ef = std::max(200, static_cast(topk) * 2); + search_params.num_threads = 1; + + hnsw::search(dev_resources, + search_params, + hnsw_index, + queries, + indices_hnsw_host.view(), + distances_hnsw_host.view()); + + auto neighbors = raft::make_device_matrix(dev_resources, n_queries, topk); + auto distances = raft::make_device_matrix(dev_resources, n_queries, topk); + auto neighbors_host = raft::make_host_matrix(n_queries, topk); + for (int64_t i = 0; i < n_queries; ++i) { + for (int64_t j = 0; j < topk; ++j) { + neighbors_host(i, j) = static_cast(indices_hnsw_host(i, j)); + } + } + + raft::copy(neighbors.data_handle(), + neighbors_host.data_handle(), + n_queries * topk, + raft::resource::get_cuda_stream(dev_resources)); + raft::copy(distances.data_handle(), + distances_hnsw_host.data_handle(), + n_queries * topk, + raft::resource::get_cuda_stream(dev_resources)); + raft::resource::sync_stream(dev_resources); + + print_results(dev_resources, neighbors.view(), distances.view()); +} + +} // namespace + +int main() +{ + raft::device_resources dev_resources; + + rmm::mr::pool_memory_resource pool_mr(rmm::mr::get_current_device_resource_ref(), + 1024 * 1024 * 1024ull); + rmm::mr::set_current_device_resource(pool_mr); + +#if HNSW_ACE_LAYERED_USE_QUANTIZATION + std::cout << "[stage 1] Generate and quantize dataset (float -> int8)" << std::endl; +#else + std::cout << "[stage 1] Generate dataset (float)" << std::endl; +#endif + + int64_t n_samples = 10000; + int64_t n_dim = 90; + int64_t n_queries = 10; + auto dataset = raft::make_device_matrix(dev_resources, n_samples, n_dim); + auto queries = raft::make_device_matrix(dev_resources, n_queries, n_dim); + generate_dataset(dev_resources, dataset.view(), queries.view()); + + auto dataset_host = raft::make_host_matrix(n_samples, n_dim); + auto queries_host = raft::make_host_matrix(n_queries, n_dim); + raft::copy(dataset_host.data_handle(), + dataset.data_handle(), + dataset.extent(0) * dataset.extent(1), + raft::resource::get_cuda_stream(dev_resources)); + raft::copy(queries_host.data_handle(), + queries.data_handle(), + queries.extent(0) * queries.extent(1), + raft::resource::get_cuda_stream(dev_resources)); + raft::resource::sync_stream(dev_resources); + + auto dataset_host_view = raft::make_host_matrix_view( + dataset_host.data_handle(), n_samples, n_dim); + auto queries_host_view = raft::make_host_matrix_view( + queries_host.data_handle(), n_queries, n_dim); + + std::filesystem::create_directories(kBuildDir); + +#if HNSW_ACE_LAYERED_USE_QUANTIZATION + auto dataset_i8 = quantize_dataset(dev_resources, dataset_host_view); + auto dataset_i8_view = raft::make_host_matrix_view( + dataset_i8.data_handle(), n_samples, n_dim); + auto dataset_path = + write_local_dataset(dataset_host_view, std::string{kBuildDir} + "/dataset.npy"); + auto hnsw_params = make_hnsw_ace_params(kBuildDir); + + std::cout << "[stage 2] Build graph-only HNSW artifact from int8 data with ACE" << std::endl; + auto artifact_path = hnsw_build(dev_resources, hnsw_params, dataset_i8_view); + + std::cout << "[stage 3] Attach original float dataset" << std::endl; + auto hnsw_index = hnsw_deserialize(dev_resources, artifact_path, dataset_path); + + std::cout << "[stage 4] Search float HNSW index" << std::endl; + hnsw_search(dev_resources, *hnsw_index, queries_host_view); +#else + auto dataset_path = + write_local_dataset(dataset_host_view, std::string{kBuildDir} + "/dataset.npy"); + auto hnsw_params = make_hnsw_ace_params(kBuildDir); + + std::cout << "[stage 2] Build layered HNSW index with ACE" << std::endl; + auto artifact_path = hnsw_build(dev_resources, hnsw_params, dataset_host_view); + + std::cout << "[stage 3] Deserialize layered HNSW index" << std::endl; + auto hnsw_index = hnsw_deserialize(dev_resources, artifact_path, dataset_path); + + std::cout << "[stage 4] Search HNSW index" << std::endl; + hnsw_search(dev_resources, *hnsw_index, queries_host_view); +#endif + + return 0; +} diff --git a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md index c7cda26d07..cdf32501d4 100644 --- a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md +++ b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md @@ -13,11 +13,13 @@ _Source header: `cuvs/neighbors/hnsw.hpp`_ Hierarchy for HNSW index when converting from CAGRA index -NOTE: When the value is `NONE`, the HNSW index is built as a base-layer-only index. +NOTE: When the value is `NONE`, the HNSW index is built as a base-layer-only index. When the value is `CPU`, a full index is built with a CPU-constructed hierarchy. When the value is `GPU`, a full index is built with a GPU-constructed hierarchy. ```cpp enum class HnswHierarchy { - NONE + NONE, + CPU, + GPU }; ``` @@ -26,6 +28,54 @@ enum class HnswHierarchy { | Name | Value | | --- | --- | | `NONE` | `` | +| `CPU` | `` | +| `GPU` | `` | + + +### neighbors::hnsw::HnswOutputFormat + +Output artifact format for an HNSW index + +`HNSWLIB` produces the standard hnswlib index format. `GRAPH_ONLY` stores the graph separately from vectors. The current implementation requires `HnswHierarchy::GPU` and disk-backed ACE. Load a graph-only artifact with the two-filename `deserialize` overload, which reads the vectors from a separate local dataset. + +```cpp +enum class HnswOutputFormat { + HNSWLIB, + GRAPH_ONLY +}; +``` + +**Values** + +| Name | Value | +| --- | --- | +| `HNSWLIB` | `` | +| `GRAPH_ONLY` | `` | + + +### neighbors::hnsw::index_params + +```cpp +struct index_params : cuvs::neighbors::index_params { + HnswHierarchy hierarchy; + HnswOutputFormat output_format; + int ef_construction; + int num_threads; + size_t M; + std::variant graph_build_params; +}; +``` + +**Fields** + +| Name | Type | Description | +| --- | --- | --- | +| `hierarchy` | [`HnswHierarchy`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-hnswhierarchy) | Hierarchy build type for HNSW index when converting from CAGRA index | +| `output_format` | [`HnswOutputFormat`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-hnswoutputformat) | Output artifact format. Graph-only output currently requires a GPU hierarchy and disk-backed ACE. | +| `ef_construction` | `int` | Size of the candidate list during hierarchy construction when hierarchy is `CPU` | +| `num_threads` | `int` | Number of host threads to use to construct hierarchy when hierarchy is `CPU` or `GPU`. When the value is 0, the number of threads is automatically determined to the maximum number of threads available.
NOTE: When hierarchy is `GPU`, while the majority of the work is done on the GPU, initialization of the HNSW index itself and some other work is parallelized with the help of CPU threads. | +| `M` | `size_t` | HNSW M parameter: number of bi-directional links per node (used when building with ACE). | +| `graph_build_params` | `std::variant` | Parameters to fine tune GPU graph building. By default we select the parameters based on dataset shape and HNSW build parameters. You can override these parameters to fine tune the graph building process as described in the CAGRA build docs.

Set ace_params to configure ACE (Augmented Core Extraction) parameters for building a GPU-accelerated HNSW index. ACE enables building indexes for datasets too large to fit in GPU memory.

When ACE writes to disk, `build_dir` may already exist, but ACE's named CAGRA artifacts and the selected HNSW output (`hnsw_index.bin` or `hnsw_index.cuvs`) must not already exist. Simultaneous builds must use different directories. The HNSW output is published only after it is fully serialized; however, the complete build is not transactional: if HNSW conversion fails after CAGRA succeeds, the completed CAGRA artifacts remain in the directory. | ### neighbors::hnsw::[[deprecated @@ -51,7 +101,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | -| `arg1` | | `"Use cagra::index_params::from_hnsw_params instead"` | | +| `arg1` | | [`"Use cagra::index_params::from_hnsw_params instead"`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | | **Returns** @@ -75,7 +125,10 @@ struct index; load a base-layer-only hnswlib index originally saved from a built CAGRA index. This is a virtual class and it cannot be used directly. To create an index, use the factory function `cuvs::neighbors::hnsw::from_cagra` from the header `cuvs/neighbors/hnsw.hpp` ```cpp -index(int dim, cuvs::distance::DistanceType metric, HnswHierarchy hierarchy = HnswHierarchy::NONE); +index(int dim, +cuvs::distance::DistanceType metric, +HnswHierarchy hierarchy = HnswHierarchy::NONE, +HnswOutputFormat output_format = HnswOutputFormat::HNSWLIB); ``` **Parameters** @@ -85,6 +138,7 @@ index(int dim, cuvs::distance::DistanceType metric, HnswHierarchy hierarchy = Hn | `dim` | in | `int` | dimensions of the training dataset | | `metric` | in | [`cuvs::distance::DistanceType`](/api-reference/cpp-api-distance-distance#distance-distancetype) | distance metric to search. Supported metrics ("L2Expanded", "InnerProduct") | | `hierarchy` | in | [`HnswHierarchy`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-hnswhierarchy) | hierarchy used for upper HNSW layers
Default: `HnswHierarchy::NONE`. | +| `output_format` | in | [`HnswOutputFormat`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-hnswoutputformat) | output artifact format
Default: `HnswOutputFormat::HNSWLIB`. | **Returns** @@ -181,7 +235,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters including ACE configuration | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters including ACE configuration | | `dataset` | in | `raft::host_matrix_view` | a host matrix view to a row-major matrix [n_rows, dim] | **Returns** @@ -212,7 +266,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters including ACE configuration | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters including ACE configuration | | `dataset` | in | `raft::host_matrix_view` | a host matrix view to a row-major matrix [n_rows, dim] | **Returns** @@ -243,7 +297,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters including ACE configuration | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters including ACE configuration | | `dataset` | in | `raft::host_matrix_view` | a host matrix view to a row-major matrix [n_rows, dim] | **Returns** @@ -274,7 +328,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters including ACE configuration | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters including ACE configuration | | `dataset` | in | `raft::host_matrix_view` | a host matrix view to a row-major matrix [n_rows, dim] | **Returns** @@ -290,6 +344,7 @@ Construct an hnswlib index from a CAGRA index NOTE: When `hnsw::index_params.hie 1. `NONE`: This method uses the filesystem to write the CAGRA index in `/tmp/<random_number>.bin` before reading it as an hnswlib index, then deleting the temporary file. The returned index is immutable and can only be searched by the hnswlib wrapper in cuVS, as the format is not compatible with the original hnswlib. 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The serialized index is also compatible with the original hnswlib library. +3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset. ```cpp std::unique_ptr> from_cagra( @@ -307,7 +362,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | | `cagra_index` | in | `const cuvs::neighbors::cagra::device_padded_index&` | cagra index | | `dataset` | in | `std::optional>` | optional dataset to avoid extra memory copy when hierarchy is `CPU`

Default: `std::nullopt`. | @@ -321,6 +376,7 @@ Construct an hnswlib index from a CAGRA index NOTE: When `hnsw::index_params.hie 1. `NONE`: This method uses the filesystem to write the CAGRA index in `/tmp/<random_number>.bin` before reading it as an hnswlib index, then deleting the temporary file. The returned index is immutable and can only be searched by the hnswlib wrapper in cuVS, as the format is not compatible with the original hnswlib. 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The serialized index is also compatible with the original hnswlib library. +3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset. ```cpp std::unique_ptr> from_cagra( @@ -338,7 +394,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | | `cagra_index` | in | `const cuvs::neighbors::cagra::device_padded_index&` | cagra index | | `dataset` | in | `std::optional>` | optional dataset to avoid extra memory copy when hierarchy is `CPU`

Default: `std::nullopt`. | @@ -352,6 +408,7 @@ Construct an hnswlib index from a CAGRA index NOTE: When `hnsw::index_params.hie 1. `NONE`: This method uses the filesystem to write the CAGRA index in `/tmp/<random_number>.bin` before reading it as an hnswlib index, then deleting the temporary file. The returned index is immutable and can only be searched by the hnswlib wrapper in cuVS, as the format is not compatible with the original hnswlib. 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The serialized index is also compatible with the original hnswlib library. +3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset. ```cpp std::unique_ptr> from_cagra( @@ -369,7 +426,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | | `cagra_index` | in | `const cuvs::neighbors::cagra::device_padded_index&` | cagra index | | `dataset` | in | `std::optional>` | optional dataset to avoid extra memory copy when hierarchy is `CPU`

Default: `std::nullopt`. | @@ -383,6 +440,7 @@ Construct an hnswlib index from a CAGRA index NOTE: When `hnsw::index_params.hie 1. `NONE`: This method uses the filesystem to write the CAGRA index in `/tmp/<random_number>.bin` before reading it as an hnswlib index, then deleting the temporary file. The returned index is immutable and can only be searched by the hnswlib wrapper in cuVS, as the format is not compatible with the original hnswlib. 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The serialized index is also compatible with the original hnswlib library. +3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset. ```cpp std::unique_ptr> from_cagra( @@ -400,7 +458,7 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | | `cagra_index` | in | `const cuvs::neighbors::cagra::device_padded_index&` | cagra index | | `dataset` | in | `std::optional>` | optional dataset to avoid extra memory copy when hierarchy is `CPU`

Default: `std::nullopt`. | @@ -428,7 +486,7 @@ When the index has an attached device dataset view, `dataset` may be omitted. Ot | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | | `raft::resources const&` | | -| `params` | | `const index_params&` | | +| `params` | | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | | | `cagra_index` | | `const cuvs::neighbors::cagra::device_standard_index&` | | | `dataset` | | `std::optional>` | Default: `std::nullopt`. | @@ -454,7 +512,7 @@ std::nullopt); | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | | `raft::resources const&` | | -| `params` | | `const index_params&` | | +| `params` | | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | | | `cagra_index` | | `const cuvs::neighbors::cagra::host_padded_index&` | | | `dataset` | | `std::optional>` | Default: `std::nullopt`. | @@ -480,7 +538,7 @@ std::nullopt); | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | | `raft::resources const&` | | -| `params` | | `const index_params&` | | +| `params` | | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | | | `cagra_index` | | `const cuvs::neighbors::cagra::host_standard_index&` | | | `dataset` | | `std::optional>` | Default: `std::nullopt`. | @@ -621,7 +679,7 @@ struct search_params : cuvs::neighbors::search_params { ### neighbors::hnsw::search -Search HNSW index constructed from a CAGRA index NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is `NONE`, as the format is not compatible with the original hnswlib. +Search HNSW index constructed from a CAGRA index NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is `NONE`, as the format is not compatible with the original hnswlib. When `output_format` is `GRAPH_ONLY`, search uses the in-memory index reconstructed from the graph-only artifact produced by the two-filename `deserialize` overload. ```cpp void search(raft::resources const& res, @@ -640,7 +698,7 @@ Usage example: | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | | `params` | in | [`const search_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-search-params) | configure the search | -| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | cagra index | +| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | HNSW index | | `queries` | in | `raft::host_matrix_view` | a host matrix view to a row-major matrix [n_queries, index->dim()] | | `neighbors` | out | `raft::host_matrix_view` | a host matrix view to the indices of the neighbors in the source dataset [n_queries, k] | | `distances` | out | `raft::host_matrix_view` | a host matrix view to the distances to the selected neighbors [n_queries, k] | @@ -651,7 +709,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::search` -Search HNSW index constructed from a CAGRA index NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is `NONE`, as the format is not compatible with the original hnswlib. +Search HNSW index constructed from a CAGRA index NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is `NONE`, as the format is not compatible with the original hnswlib. When `output_format` is `GRAPH_ONLY`, search uses the in-memory index reconstructed from the graph-only artifact produced by the two-filename `deserialize` overload. ```cpp void search(raft::resources const& res, @@ -670,7 +728,7 @@ Usage example: | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | | `params` | in | [`const search_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-search-params) | configure the search | -| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | cagra index | +| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | HNSW index | | `queries` | in | `raft::host_matrix_view` | a host matrix view to a row-major matrix [n_queries, index->dim()] | | `neighbors` | out | `raft::host_matrix_view` | a host matrix view to the indices of the neighbors in the source dataset [n_queries, k] | | `distances` | out | `raft::host_matrix_view` | a host matrix view to the distances to the selected neighbors [n_queries, k] | @@ -681,7 +739,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::search` -Search HNSWindex constructed from a CAGRA index NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is `NONE`, as the format is not compatible with the original hnswlib. +Search HNSWindex constructed from a CAGRA index NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is `NONE`, as the format is not compatible with the original hnswlib. When `output_format` is `GRAPH_ONLY`, search uses the in-memory index reconstructed from the graph-only artifact produced by the two-filename `deserialize` overload. ```cpp void search(raft::resources const& res, @@ -700,7 +758,7 @@ Usage example: | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | | `params` | in | [`const search_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-search-params) | configure the search | -| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | cagra index | +| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | HNSW index | | `queries` | in | `raft::host_matrix_view` | a host matrix view to a row-major matrix [n_queries, index->dim()] | | `neighbors` | out | `raft::host_matrix_view` | a host matrix view to the indices of the neighbors in the source dataset [n_queries, k] | | `distances` | out | `raft::host_matrix_view` | a host matrix view to the distances to the selected neighbors [n_queries, k] | @@ -711,7 +769,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::search` -Search HNSW index constructed from a CAGRA index NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is `NONE`, as the format is not compatible with the original hnswlib. +Search HNSW index constructed from a CAGRA index NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS when the hierarchy is `NONE`, as the format is not compatible with the original hnswlib. When `output_format` is `GRAPH_ONLY`, search uses the in-memory index reconstructed from the graph-only artifact produced by the two-filename `deserialize` overload. ```cpp void search(raft::resources const& res, @@ -730,7 +788,7 @@ Usage example: | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | | `params` | in | [`const search_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-search-params) | configure the search | -| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | cagra index | +| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | HNSW index | | `queries` | in | `raft::host_matrix_view` | a host matrix view to a row-major matrix [n_queries, index->dim()] | | `neighbors` | out | `raft::host_matrix_view` | a host matrix view to the indices of the neighbors in the source dataset [n_queries, k] | | `distances` | out | `raft::host_matrix_view` | a host matrix view to the distances to the selected neighbors [n_queries, k] | @@ -739,12 +797,12 @@ Usage example: `void` -## Deserialize CAGRA index as hnswlib index +## Serialize and deserialize HNSW indexes ### neighbors::hnsw::serialize -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -757,8 +815,8 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `filename` | in | `const std::string&` | path to the file to save the serialized CAGRA index | -| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | cagra index | +| `filename` | in | `const std::string&` | path to the serialized HNSW output | +| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | HNSW index | **Returns** @@ -766,7 +824,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::serialize` -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -779,8 +837,8 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `filename` | in | `const std::string&` | path to the file to save the serialized CAGRA index | -| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | cagra index | +| `filename` | in | `const std::string&` | path to the serialized HNSW output | +| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | HNSW index | **Returns** @@ -788,7 +846,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::serialize` -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -801,8 +859,8 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `filename` | in | `const std::string&` | path to the file to save the serialized CAGRA index | -| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | cagra index | +| `filename` | in | `const std::string&` | path to the serialized HNSW output | +| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | HNSW index | **Returns** @@ -810,7 +868,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::serialize` -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -823,8 +881,8 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `filename` | in | `const std::string&` | path to the file to save the serialized CAGRA index | -| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | cagra index | +| `filename` | in | `const std::string&` | path to the serialized HNSW output | +| `idx` | in | [`const index&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | HNSW index | **Returns** @@ -833,7 +891,7 @@ Usage example: ### neighbors::hnsw::deserialize -De-serialize a CAGRA index saved to a file as an hnswlib index NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. +Deserialize an HNSWLIB index NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. This overload loads `HNSWLIB` artifacts. Use the two-filename overload for graph-only artifacts. ```cpp void deserialize(raft::resources const& res, @@ -851,8 +909,8 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters | -| `filename` | in | `const std::string&` | path to the file containing the serialized CAGRA index | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | +| `filename` | in | `const std::string&` | path to the HNSWLIB artifact | | `dim` | in | `int` | dimensions of the training dataset | | `metric` | in | [`cuvs::distance::DistanceType`](/api-reference/cpp-api-distance-distance#distance-distancetype) | distance metric to search. Supported metrics ("L2Expanded", "InnerProduct") | | `index` | out | [`index**`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | hnsw index | @@ -863,7 +921,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::deserialize` -De-serialize a CAGRA index saved to a file as an hnswlib index NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. +Deserialize an HNSWLIB index NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. This overload loads `HNSWLIB` artifacts. Use the two-filename overload for graph-only artifacts. ```cpp void deserialize(raft::resources const& res, @@ -881,8 +939,8 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters | -| `filename` | in | `const std::string&` | path to the file containing the serialized CAGRA index | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | +| `filename` | in | `const std::string&` | path to the HNSWLIB artifact | | `dim` | in | `int` | dimensions of the training dataset | | `metric` | in | [`cuvs::distance::DistanceType`](/api-reference/cpp-api-distance-distance#distance-distancetype) | distance metric to search. Supported metrics ("L2Expanded", "InnerProduct") | | `index` | out | [`index**`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | hnsw index | @@ -893,7 +951,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::deserialize` -De-serialize a CAGRA index saved to a file as an hnswlib index NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. +Deserialize an HNSWLIB index NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. This overload loads `HNSWLIB` artifacts. Use the two-filename overload for graph-only artifacts. ```cpp void deserialize(raft::resources const& res, @@ -911,8 +969,8 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters | -| `filename` | in | `const std::string&` | path to the file containing the serialized CAGRA index | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | +| `filename` | in | `const std::string&` | path to the HNSWLIB artifact | | `dim` | in | `int` | dimensions of the training dataset | | `metric` | in | [`cuvs::distance::DistanceType`](/api-reference/cpp-api-distance-distance#distance-distancetype) | distance metric to search. Supported metrics ("L2Expanded", "InnerProduct") | | `index` | out | [`index**`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | hnsw index | @@ -923,7 +981,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::deserialize` -De-serialize a CAGRA index saved to a file as an hnswlib index NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. +Deserialize an HNSWLIB index NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. This overload loads `HNSWLIB` artifacts. Use the two-filename overload for graph-only artifacts. ```cpp void deserialize(raft::resources const& res, @@ -941,8 +999,8 @@ Usage example: | Name | Direction | Type | Description | | --- | --- | --- | --- | | `res` | in | `raft::resources const&` | raft resources | -| `params` | in | `const index_params&` | hnsw index parameters | -| `filename` | in | `const std::string&` | path to the file containing the serialized CAGRA index | +| `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | +| `filename` | in | `const std::string&` | path to the HNSWLIB artifact | | `dim` | in | `int` | dimensions of the training dataset | | `metric` | in | [`cuvs::distance::DistanceType`](/api-reference/cpp-api-distance-distance#distance-distancetype) | distance metric to search. Supported metrics ("L2Expanded", "InnerProduct") | | `index` | out | [`index**`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | hnsw index | @@ -950,3 +1008,97 @@ Usage example: **Returns** `void` + +**Additional overload:** `neighbors::hnsw::deserialize` + +Deserialize a graph-only HNSW artifact and attach its dataset + +```cpp +void deserialize(raft::resources const& res, +const std::string& graph_filename, +const std::string& dataset_filename, +index** index); +``` + +The graph artifact supplies the index dimensions, metric, and construction metadata. The attached dataset must have the recorded shape, but its element type may differ from the type used to construct the graph. The output pointer selects the attached dataset type. The loader accepts row-major `.npy` files and ANN benchmark binary files with a `[uint32 rows, uint32 cols]` header. Binary extensions must match the output index type: `.fbin` for `float`, `.f16bin` or `.fp16.fbin` for `half`, `.u8bin` for `uint8_t`, and `.i8bin` for `int8_t`. + +Usage example: + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `res` | in | `raft::resources const&` | raft resources | +| `graph_filename` | in | `const std::string&` | path to the graph-only HNSW artifact | +| `dataset_filename` | in | `const std::string&` | path to the local dataset | +| `index` | out | [`index**`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | reconstructed HNSW index | + +**Returns** + +`void` + +**Additional overload:** `neighbors::hnsw::deserialize` + +```cpp +void deserialize(raft::resources const& res, +const std::string& graph_filename, +const std::string& dataset_filename, +index** index); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `res` | | `raft::resources const&` | | +| `graph_filename` | | `const std::string&` | | +| `dataset_filename` | | `const std::string&` | | +| `index` | | [`index**`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | | + +**Returns** + +`void` + +**Additional overload:** `neighbors::hnsw::deserialize` + +```cpp +void deserialize(raft::resources const& res, +const std::string& graph_filename, +const std::string& dataset_filename, +index** index); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `res` | | `raft::resources const&` | | +| `graph_filename` | | `const std::string&` | | +| `dataset_filename` | | `const std::string&` | | +| `index` | | [`index**`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | | + +**Returns** + +`void` + +**Additional overload:** `neighbors::hnsw::deserialize` + +```cpp +void deserialize(raft::resources const& res, +const std::string& graph_filename, +const std::string& dataset_filename, +index** index); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `res` | | `raft::resources const&` | | +| `graph_filename` | | `const std::string&` | | +| `dataset_filename` | | `const std::string&` | | +| `index` | | [`index**`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index) | | + +**Returns** + +`void`