diff --git a/c/include/cuvs/neighbors/hnsw.h b/c/include/cuvs/neighbors/hnsw.h index 58ea5022e0..442a793e91 100644 --- a/c/include/cuvs/neighbors/hnsw.h +++ b/c/include/cuvs/neighbors/hnsw.h @@ -594,6 +594,81 @@ CUVS_EXPORT cuvsError_t cuvsHnswDeserialize(cuvsResources_t res, * @} */ +/** + * @defgroup hnsw_c_index_materialize Materialize a layered HNSW artifact to an hnswlib index + * @{ + */ + +/** + * @brief Parameters for materializing a layered HNSW artifact into an hnswlib index on disk. + */ +struct cuvsHnswMaterializeParams { + /** + * Local dataset path holding the original-ID-ordered vectors used to build the artifact. + * + * Supported formats match layered deserialization: `.npy` and ANN benchmark `*.bin` files with a + * `[uint32 rows, uint32 cols]` header (`.fbin`, `.f16bin`, `.u8bin`, `.i8bin`). + */ + const char* dataset_path; + /** + * Upper bound on host memory (in GiB) used for the base-topology reorder buffer. + * + * When `<= 0`, the whole base topology is reordered in a single in-memory pass (no temporary + * files). When set, the base topology is reordered through bucketed temporary files so that peak + * host memory stays close to this budget. + */ + double max_host_memory_gb; + /** Number of host threads to use. When `0`, the maximum number of threads is used. */ + int num_threads; +}; + +typedef struct cuvsHnswMaterializeParams* cuvsHnswMaterializeParams_t; + +/** + * @brief Allocate HNSW materialize params, and populate with default values + * + * @param[in] params cuvsHnswMaterializeParams_t to allocate + * @return cuvsError_t + */ +CUVS_EXPORT cuvsError_t cuvsHnswMaterializeParamsCreate(cuvsHnswMaterializeParams_t* params); + +/** + * @brief De-allocate HNSW materialize params + * + * @param[in] params cuvsHnswMaterializeParams_t to de-allocate + * @return cuvsError_t + */ +CUVS_EXPORT cuvsError_t cuvsHnswMaterializeParamsDestroy(cuvsHnswMaterializeParams_t params); + +/** + * @brief Materialize a layered HNSW artifact into a standard hnswlib index file on disk. + * + * Materializes a `GRAPH_ONLY` artifact (graph topology only, stored in ACE order) plus a + * local dataset into a standard hnswlib index file, without ever holding the full materialized + * index in host memory. The resulting file is compatible with the original hnswlib library and can + * be read back through `cuvsHnswDeserialize` with `hierarchy == CPU`. The element data type + * (`float`, `half`, `uint8_t` or `int8_t`) is inferred from the external dataset. GRAPH_ONLY + * artifacts are currently produced through the C++ API. + * + * @param[in] res cuvsResources_t opaque C handle + * @param[in] params cuvsHnswMaterializeParams_t materialization parameters + * @param[in] layered_artifact_path path to the layered HNSW artifact + * @param[in] output_path path to the hnswlib index file to write + * @param[in] dim the dimension of the vectors in the index + * @param[in] metric the distance metric used to build the index + * @return cuvsError_t + */ +CUVS_EXPORT cuvsError_t cuvsHnswMaterializeToHnswlib(cuvsResources_t res, + cuvsHnswMaterializeParams_t params, + const char* layered_artifact_path, + const char* output_path, + int dim, + cuvsDistanceType metric); + +/** + * @} + */ + #ifdef __cplusplus } #endif diff --git a/c/src/neighbors/hnsw.cpp b/c/src/neighbors/hnsw.cpp index 1f19f33875..0dbdd0c4c7 100644 --- a/c/src/neighbors/hnsw.cpp +++ b/c/src/neighbors/hnsw.cpp @@ -408,3 +408,41 @@ extern "C" cuvsError_t cuvsHnswDeserialize(cuvsResources_t res, } }); } + +extern "C" cuvsError_t cuvsHnswMaterializeParamsCreate(cuvsHnswMaterializeParams_t* params) +{ + return cuvs::core::translate_exceptions([=] { + *params = new cuvsHnswMaterializeParams{ + .dataset_path = nullptr, .max_host_memory_gb = 0, .num_threads = 0}; + }); +} + +extern "C" cuvsError_t cuvsHnswMaterializeParamsDestroy(cuvsHnswMaterializeParams_t params) +{ + return cuvs::core::translate_exceptions([=] { delete params; }); +} + +extern "C" cuvsError_t cuvsHnswMaterializeToHnswlib(cuvsResources_t res, + cuvsHnswMaterializeParams_t params, + const char* layered_artifact_path, + const char* output_path, + int dim, + cuvsDistanceType metric) +{ + return cuvs::core::translate_exceptions([=] { + auto res_ptr = reinterpret_cast(res); + auto cpp_params = cuvs::neighbors::hnsw::materialize_params(); + if (params->dataset_path != nullptr) { + cpp_params.dataset_path = std::string(params->dataset_path); + } + cpp_params.max_host_memory_gb = params->max_host_memory_gb; + cpp_params.num_threads = params->num_threads; + auto metric_type = static_cast(metric); + cuvs::neighbors::hnsw::materialize_to_hnswlib(*res_ptr, + cpp_params, + std::string(layered_artifact_path), + std::string(output_path), + dim, + metric_type); + }); +} diff --git a/cpp/include/cuvs/neighbors/hnsw.hpp b/cpp/include/cuvs/neighbors/hnsw.hpp index fe412e8cc1..960aaf55ee 100644 --- a/cpp/include/cuvs/neighbors/hnsw.hpp +++ b/cpp/include/cuvs/neighbors/hnsw.hpp @@ -1451,6 +1451,76 @@ void deserialize(raft::resources const& res, * @} */ +/** + * @defgroup hnsw_cpp_index_materialize Materialize a layered HNSW artifact into an hnswlib index + * @{ + */ + +/** + * @brief Parameters for materializing a layered HNSW artifact into an hnswlib index on disk. + */ +struct materialize_params { + /** Local dataset path holding the original-ID-ordered vectors used to build the artifact. + * + * Supported formats match layered deserialization: `.npy` and ANN benchmark `*.bin` files with a + * `[uint32 rows, uint32 cols]` header (`.fbin`, `.f16bin`, `.u8bin`, `.i8bin`). + */ + std::string dataset_path; + + /** Upper bound on host memory (in GiB) used for the base-topology reorder buffer. + * + * When `<= 0`, the whole base topology is reordered in a single in-memory pass (no temporary + * files). When set, the base topology is reordered through bucketed temporary files so that + * peak host memory stays close to this budget, at the cost of writing and re-reading the + * (small) base-topology section once. + */ + double max_host_memory_gb = 0; + + /** Number of host threads to use. When `0`, the maximum number of threads is used. */ + int num_threads = 0; +}; + +/** + * @brief Materialize a layered HNSW artifact into a standard hnswlib index file on disk. + * + * Materializes a `GRAPH_ONLY` artifact (graph topology only, stored in ACE order) plus a + * local dataset into a standard hnswlib index file, without ever holding the full materialized + * index in host memory. The materialization reorders the base topology from ACE order to + * original-id order and interleaves the vectors, emitting the output with sequential disk IO. The + * resulting file is compatible with the original hnswlib library (`loadIndex`) and can be read back + * through `cuvs::neighbors::hnsw::deserialize` with `hierarchy == HnswHierarchy::CPU`. + * + * The element data type (`float`, `half`, `uint8_t` or `int8_t`) is inferred from the external + * dataset, so materialization supports an original dataset dtype that differs from the graph's + * construction dtype. + * + * @param[in] res raft resources + * @param[in] params materialization parameters (dataset path, host-memory budget, threads) + * @param[in] layered_artifact_path path to the layered HNSW artifact + * @param[in] output_path path to the hnswlib index file to write + * @param[in] dim dimensions of the training dataset + * @param[in] metric distance metric. Supported metrics ("L2Expanded", "InnerProduct") + * + * Usage example: + * @code{.cpp} + * using namespace cuvs::neighbors; + * hnsw::materialize_params materialize_params; + * materialize_params.dataset_path = "dataset.fbin"; + * hnsw::materialize_to_hnswlib( + * res, materialize_params, "layered_artifact.cuvs", "index.bin", dim, metric); + * @endcode + */ +void materialize_to_hnswlib(raft::resources const& res, + const materialize_params& params, + const std::string& layered_artifact_path, + const std::string& output_path, + int dim, + cuvs::distance::DistanceType metric); + +/** + * @} + */ + } // namespace hnsw } // namespace neighbors } // namespace CUVS_EXPORT cuvs diff --git a/cpp/include/cuvs/util/file_io.hpp b/cpp/include/cuvs/util/file_io.hpp index e2e9d34026..be532726c1 100644 --- a/cpp/include/cuvs/util/file_io.hpp +++ b/cpp/include/cuvs/util/file_io.hpp @@ -443,6 +443,19 @@ void write_large_file(const file_descriptor& fd, const size_t total_bytes, const uint64_t file_offset); +/** + * @brief Pre-size a file to `total_bytes` bytes. + * + * Prefers posix_fallocate (reserves blocks up-front, avoids later ENOSPC and fragmentation), but + * falls back to ftruncate on filesystems that do not support preallocation (tmpfs and some + * NFS/overlay mounts return EOPNOTSUPP/EINVAL/ENOSYS) so the operation still succeeds there. A + * `total_bytes` of 0 is a no-op. Throws on failure (uses the descriptor's path in the message). + * + * @param fd File descriptor to pre-size + * @param total_bytes Target file size in bytes + */ +void preallocate_file(const file_descriptor& fd, const size_t total_bytes); + /** * @brief Sequential std::ostream backed by kvikio. * diff --git a/cpp/src/neighbors/detail/hnsw.hpp b/cpp/src/neighbors/detail/hnsw.hpp index 19c0afd820..e365265b7b 100644 --- a/cpp/src/neighbors/detail/hnsw.hpp +++ b/cpp/src/neighbors/detail/hnsw.hpp @@ -35,6 +35,7 @@ #include #include #include +#include #include #include #include @@ -103,6 +104,18 @@ class exclusive_hnsw_temp_file { } } + void publish_replace() + { + if (::rename(temporary_path_.c_str(), output_path_.c_str()) != 0) { + const int error = errno; + RAFT_FAIL("Cannot publish HNSW index %s (errno: %d, %s)", + output_path_.c_str(), + error, + strerror(error)); + } + temporary_path_.clear(); + } + private: void cleanup() noexcept { @@ -893,6 +906,33 @@ inline auto open_layered_dataset_file(const std::string& path) -> npy_file false}; } +inline auto layered_dataset_dtype(const std::string& path) -> layered_hnsw_dtype +{ + if (std::filesystem::path(path).extension() == ".npy") { + const auto file = open_npy_file(path); + if (file.dtype == cuvs::util::detail::numpy_dtype_string()) { + return layered_hnsw_dtype::float32; + } + if (file.dtype == cuvs::util::detail::numpy_dtype_string()) { + return layered_hnsw_dtype::float16; + } + if (file.dtype == cuvs::util::detail::numpy_dtype_string()) { + return layered_hnsw_dtype::uint8; + } + if (file.dtype == cuvs::util::detail::numpy_dtype_string()) { + return layered_hnsw_dtype::int8; + } + return layered_hnsw_dtype::unknown; + } + + const auto ext = std::filesystem::path(path).extension().string(); + if (ext == ".fbin" && !ends_with(path, ".fp16.fbin")) { return layered_hnsw_dtype::float32; } + if (ext == ".f16bin" || ends_with(path, ".fp16.fbin")) { return layered_hnsw_dtype::float16; } + if (ext == ".u8bin") { return layered_hnsw_dtype::uint8; } + if (ext == ".i8bin") { return layered_hnsw_dtype::int8; } + return layered_hnsw_dtype::unknown; +} + template void write_layered_base_links_from_disk(const CagraIndexT& index_, const cuvs::util::file_descriptor& output_fd, @@ -1177,11 +1217,7 @@ auto serialize_to_layered_hnswlib_from_disk( 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::preallocate_file(artifact_fd, final_file_size); cuvs::util::write_large_file(artifact_fd, &header, sizeof(header), 0); if (descriptors_bytes > 0) { cuvs::util::write_large_file( @@ -1323,6 +1359,45 @@ auto serialize_to_layered_hnswlib_from_disk( return artifact_file.string(); } +// Build the standard hnswlib index header (matches HierarchicalNSW::saveIndex byte layout). Single +// source of the on-disk header encoding, shared by the streaming serialize path and the +// disk-to-disk materialize path so the two cannot drift. +inline auto make_hnswlib_native_header(size_t offset_level0, + size_t n_rows, + size_t size_data_per_element, + size_t label_offset, + size_t offset_data, + int maxlevel, + int enterpoint_node, + size_t maxM, + size_t maxM0, + size_t M, + double mult, + size_t ef_construction) -> std::vector +{ + std::vector buffer; + auto append = [&buffer](const auto& value) { + const auto* bytes = reinterpret_cast(&value); + buffer.insert(buffer.end(), bytes, bytes + sizeof(value)); + }; + const size_t max_elements = n_rows; + const size_t cur_element_count = n_rows; + append(offset_level0); + append(max_elements); + append(cur_element_count); + append(size_data_per_element); + append(label_offset); + append(offset_data); + append(maxlevel); + append(enterpoint_node); + append(maxM); + append(maxM0); + append(M); + append(mult); + append(ef_construction); + return buffer; +} + // 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. @@ -1408,33 +1483,21 @@ void serialize_to_hnswlib_batched(raft::resources const& res, appr_algo->maxM0_, appr_algo->M_); - // offset_level_0 - os.write(reinterpret_cast(&appr_algo->offsetLevel0_), sizeof(std::size_t)); - // 8 max_element - override with n_rows - size_t num_elements = (size_t)n_rows; - os.write(reinterpret_cast(&num_elements), sizeof(std::size_t)); - // 16 curr_element_count - override with n_rows - os.write(reinterpret_cast(&num_elements), sizeof(std::size_t)); - // 24 size_data_per_element - os.write(reinterpret_cast(&appr_algo->size_data_per_element_), sizeof(std::size_t)); - // 32 label_offset - os.write(reinterpret_cast(&appr_algo->label_offset_), sizeof(std::size_t)); - // 40 offset_data - os.write(reinterpret_cast(&appr_algo->offsetData_), sizeof(std::size_t)); - // 48 maxlevel - os.write(reinterpret_cast(&appr_algo->maxlevel_), sizeof(int)); - // 52 enterpoint_node - os.write(reinterpret_cast(&appr_algo->enterpoint_node_), sizeof(int)); - // 56 maxM - os.write(reinterpret_cast(&appr_algo->maxM_), sizeof(std::size_t)); - // 64 maxM0 - os.write(reinterpret_cast(&appr_algo->maxM0_), sizeof(std::size_t)); - // 72 M - os.write(reinterpret_cast(&appr_algo->M_), sizeof(std::size_t)); - // 80 mult - os.write(reinterpret_cast(&appr_algo->mult_), sizeof(double)); - // 88 ef_construction - os.write(reinterpret_cast(&appr_algo->ef_construction_), sizeof(std::size_t)); + // Write the hnswlib index header (max_element / cur_element_count overridden with n_rows) using + // the shared encoder so the byte layout stays in lockstep with the materialize path. + const auto native_header = make_hnswlib_native_header(appr_algo->offsetLevel0_, + static_cast(n_rows), + appr_algo->size_data_per_element_, + appr_algo->label_offset_, + appr_algo->offsetData_, + appr_algo->maxlevel_, + appr_algo->enterpoint_node_, + appr_algo->maxM_, + appr_algo->maxM0_, + appr_algo->M_, + appr_algo->mult_, + appr_algo->ef_construction_); + os.write(native_header.data(), static_cast(native_header.size())); // host queries auto host_query_set = @@ -2512,15 +2575,34 @@ void serialize(raft::resources const& res, const std::string& filename, const in hnswlib_index->saveIndex(filename); } +// Parsed + fully validated view of a layered HNSW artifact: header metadata, the payload section +// offsets, and the per-row levels array. Both the in-memory deserialize path and the disk-to-disk +// materialize path consume this so the artifact format (offset arithmetic and the size-invariant +// checks) lives in a single place and cannot drift between them. +struct layered_artifact_view { + layered_hnsw_file_metadata metadata; + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2Expanded; + size_t artifact_size = 0; + size_t levels_offset = 0; + size_t base_nodes_offset = 0; + size_t base_links_offset = 0; + size_t upper_nodes_offset = 0; + size_t upper_links_offset = 0; + std::vector levels; // per-row top level, length n_rows +}; + +// Reads and validates the artifact header/descriptors/levels from an already-open artifact fd. +// Enforces every structural invariant (magic, version, dtype, metric, dim, layer count, section +// sizes, file length, and the levels max-level) BEFORE the caller allocates buffers sized from +// these fields, so a corrupt/truncated/crafted header cannot drive an out-of-bounds read or write. template -auto deserialize_layered_hnswlib(raft::resources const& res, - const std::string& artifact_path, - const std::string& dataset_path) -> std::unique_ptr> +auto read_and_validate_layered_artifact( + const cuvs::util::file_descriptor& artifact_fd, + const std::string& artifact_path, + std::optional expected_dim = std::nullopt, + std::optional expected_metric = std::nullopt) + -> layered_artifact_view { - 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, @@ -2533,34 +2615,49 @@ auto deserialize_layered_hnswlib(raft::resources const& res, 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, + const auto artifact_metric = static_cast(header.metric); + RAFT_EXPECTS(artifact_metric == cuvs::distance::DistanceType::L2Expanded || + artifact_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 artifact_dim = static_cast(header.dim); + if (expected_dim.has_value()) { + RAFT_EXPECTS(artifact_dim == *expected_dim, + "Layered HNSW artifact dim (%d) does not match requested dim (%d)", + artifact_dim, + *expected_dim); + } + if (expected_metric.has_value()) { + RAFT_EXPECTS(artifact_metric == *expected_metric, + "Layered HNSW artifact metric (%s) does not match requested metric (%s)", + metric_name(artifact_metric), + metric_name(*expected_metric)); + } + + layered_artifact_view view; + view.metric = artifact_metric; + view.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, + RAFT_EXPECTS(descriptors_offset + descriptors_bytes <= view.artifact_size, "Layered HNSW layer descriptors are outside artifact: offset=%zu size=%zu " "artifact=%zu", descriptors_offset, descriptors_bytes, - artifact_size); + view.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); + auto& metadata = view.metadata; + 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), @@ -2569,7 +2666,6 @@ auto deserialize_layered_hnswlib(raft::resources const& res, 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"); @@ -2583,20 +2679,69 @@ auto deserialize_layered_hnswlib(raft::resources const& res, metadata.layers.size(), metadata.maxlevel); - RAFT_EXPECTS(!dataset_path.empty(), "Layered HNSW deserialization requires a dataset filename"); + // Section-size invariants: every buffer below is sized from n_rows / upper_nodes_count, so the + // declared section byte counts must match exactly or a later read/scatter would overrun. + 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); 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, + view.levels_offset = payload_offset; + view.base_nodes_offset = view.levels_offset + metadata.levels_bytes; + view.base_links_offset = view.base_nodes_offset + metadata.base_nodes_bytes; + view.upper_nodes_offset = view.base_links_offset + metadata.base_links_bytes; + view.upper_links_offset = view.upper_nodes_offset + metadata.upper_nodes_bytes; + const auto expected_file_size = view.upper_links_offset + metadata.upper_links_bytes; + RAFT_EXPECTS(view.artifact_size >= expected_file_size, "Layered HNSW artifact is truncated: expected at least %zu bytes, got %zu", expected_file_size, - artifact_size); + view.artifact_size); + + view.levels.resize(metadata.n_rows); + cuvs::util::read_large_file( + artifact_fd, view.levels.data(), metadata.levels_bytes, view.levels_offset); + const auto max_level_in_levels = *std::max_element(view.levels.begin(), view.levels.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); + return view; +} + +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); + auto view = read_and_validate_layered_artifact(artifact_fd, artifact_path); + auto& metadata = view.metadata; + const auto dim = static_cast(metadata.dim); + const auto metric = view.metric; + const auto artifact_size = view.artifact_size; + const auto base_nodes_offset = view.base_nodes_offset; + const auto base_links_offset = view.base_links_offset; + const auto upper_nodes_offset = view.upper_nodes_offset; + const auto upper_links_offset = view.upper_links_offset; + const auto metadata_elapsed_ms = elapsed_ms_since(metadata_start_time); + + RAFT_EXPECTS(!dataset_path.empty(), "Layered HNSW deserialization requires a dataset filename"); RAFT_LOG_INFO("Layered HNSW load: metadata read in %ld ms (rows=%zu dim=%zu artifact=%.2f GiB)", metadata_elapsed_ms, metadata.n_rows, @@ -2617,21 +2762,6 @@ auto deserialize_layered_hnswlib(raft::resources const& res, 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, @@ -2659,19 +2789,8 @@ auto deserialize_layered_hnswlib(raft::resources const& res, } }; - 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); + auto& levels_u8 = view.levels; 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 = @@ -2884,6 +3003,683 @@ auto deserialize_layered_hnswlib(raft::resources const& res, return hnsw_index; } +// Disk-to-disk materialization: layered HNSW artifact -> standard hnswlib index file. +// Constants/offsets shared by the materialization helpers. +struct hnswlib_materialize_layout { + size_t n_rows = 0; + size_t dim = 0; + size_t base_link_row_bytes = 0; + size_t upper_link_row_bytes = 0; + size_t size_data_per_element = 0; + size_t offset_data = 0; + size_t label_offset = 0; + size_t data_size = 0; + // Output file offsets. + size_t base_region_offset = 0; + size_t upper_region_offset = 0; + // Artifact payload offsets. + size_t base_nodes_offset = 0; + size_t base_links_offset = 0; + size_t upper_nodes_offset = 0; + size_t upper_links_offset = 0; +}; + +class scoped_directory_cleanup { + public: + explicit scoped_directory_cleanup(std::filesystem::path path) : path_{std::move(path)} {} + scoped_directory_cleanup(const scoped_directory_cleanup&) = delete; + scoped_directory_cleanup& operator=(const scoped_directory_cleanup&) = delete; + ~scoped_directory_cleanup() + { + std::error_code ec; + std::filesystem::remove_all(path_, ec); + } + + private: + std::filesystem::path path_; +}; + +// Routes fixed-size, ID-keyed records into per-bucket temporary files and replays them per bucket. +// Used to reorder the (small) base/upper topology under a bounded host-memory budget while keeping +// all disk access sequential. +struct id_record_spiller { + std::filesystem::path dir; + size_t rows_per_bucket; + size_t num_buckets; + size_t record_bytes; + size_t buffer_cap; + std::vector fds; + std::vector offsets; + std::vector> buffers; + // One mutex per bucket so distinct buckets can be appended concurrently (add() is called from + // parallel scatter threads); same-bucket appends serialize on their own mutex. + std::vector bucket_mutexes; + + id_record_spiller(std::filesystem::path dir_, + size_t num_buckets_, + size_t rows_per_bucket_, + size_t record_bytes_, + size_t total_buffer_budget) + : dir(std::move(dir_)), + rows_per_bucket(rows_per_bucket_), + num_buckets(num_buckets_), + record_bytes(record_bytes_), + offsets(num_buckets_, 0), + buffers(num_buckets_), + bucket_mutexes(num_buckets_) + { + std::filesystem::create_directories(dir); + fds.reserve(num_buckets); + for (size_t b = 0; b < num_buckets; ++b) { + fds.emplace_back((dir / ("bucket_" + std::to_string(b) + ".tmp")).string(), + O_CREAT | O_RDWR | O_TRUNC, + 0644); + } + // Floor the per-bucket buffer so a tiny budget (many buckets) still batches many records per + // flush instead of degrading into one write syscall per record. This may exceed the nominal + // budget for pathologically small budgets, but keeps disk I/O coarse-grained and sequential. + constexpr size_t kMinRecordsPerFlush = 256; + const size_t even_share = total_buffer_budget / std::max(1, num_buckets); + buffer_cap = std::max(record_bytes * kMinRecordsPerFlush, even_share); + } + + id_record_spiller(const id_record_spiller&) = delete; + id_record_spiller& operator=(const id_record_spiller&) = delete; + id_record_spiller(id_record_spiller&&) = delete; + id_record_spiller& operator=(id_record_spiller&&) = delete; + + ~id_record_spiller() noexcept + { + fds.clear(); + std::error_code ec; + std::filesystem::remove_all(dir, ec); + } + + void flush(size_t b) + { + if (buffers[b].empty()) { return; } + cuvs::util::write_large_file(fds[b], buffers[b].data(), buffers[b].size(), offsets[b]); + offsets[b] += buffers[b].size(); + buffers[b].clear(); + } + + void add(size_t bucket, const void* record) + { + std::lock_guard guard(bucket_mutexes[bucket]); + auto& buf = buffers[bucket]; + if (buf.size() + record_bytes > buffer_cap) { flush(bucket); } + const auto* p = reinterpret_cast(record); + buf.insert(buf.end(), p, p + record_bytes); + } + + void finish_writes() + { + for (size_t b = 0; b < num_buckets; ++b) { + flush(b); + } + } + + template + void replay(size_t b, std::vector& chunk, F&& consume) + { + const size_t total = offsets[b]; + const size_t recs_per_chunk = std::max(1, (size_t{64} << 20) / record_bytes); + chunk.resize(recs_per_chunk * record_bytes); + size_t read_off = 0; + while (read_off < total) { + const size_t bytes = std::min(chunk.size(), total - read_off); + cuvs::util::read_large_file(fds[b], chunk.data(), bytes, read_off); + const size_t nrec = bytes / record_bytes; + for (size_t r = 0; r < nrec; ++r) { + consume(chunk.data() + r * record_bytes); + } + read_off += bytes; + } + } +}; + +// Streams the base topology (base_nodes + base_links) in artifact (ACE) order, invoking +// `cb(original_id, link_row_ptr)` for each row. When `parallel_threads > 1` the callback is invoked +// concurrently and must be thread-safe for distinct IDs. +template +void scatter_layered_base_links(const cuvs::util::file_descriptor& artifact_fd, + const hnswlib_materialize_layout& layout, + int parallel_threads, + F&& cb) +{ + const size_t n_rows = layout.n_rows; + const size_t row = layout.base_link_row_bytes; + const size_t batch = std::max(1, (size_t{64} << 20) / (sizeof(IdxT) + row)); + std::vector node_buf(batch); + std::vector link_buf(batch * row); + for (size_t s = 0; s < n_rows; s += batch) { + const size_t cur = std::min(batch, n_rows - s); + cuvs::util::read_large_file(artifact_fd, + node_buf.data(), + cur * sizeof(IdxT), + layout.base_nodes_offset + s * sizeof(IdxT)); + cuvs::util::read_large_file( + artifact_fd, link_buf.data(), cur * row, layout.base_links_offset + s * row); + bool invalid = false; + if (parallel_threads > 1) { +#pragma omp parallel for num_threads(parallel_threads) reduction(|| : invalid) + for (int64_t k = 0; k < static_cast(cur); ++k) { + const size_t id = static_cast(node_buf[k]); + if (id >= n_rows) { + invalid = true; + continue; + } + cb(id, link_buf.data() + static_cast(k) * row); + } + } else { + for (size_t k = 0; k < cur; ++k) { + const size_t id = static_cast(node_buf[k]); + if (id >= n_rows) { + invalid = true; + break; + } + cb(id, link_buf.data() + k * row); + } + } + RAFT_EXPECTS(!invalid, "Invalid base-layer node id in layered HNSW artifact"); + } +} + +// Streams the upper layers, invoking `cb(original_id, level, link_row_ptr)` for every promoted row. +template +void scatter_layered_upper_links(const cuvs::util::file_descriptor& artifact_fd, + const hnswlib_materialize_layout& layout, + const std::vector& layers, + const std::vector& levels_u8, + F&& cb) +{ + const size_t urow = layout.upper_link_row_bytes; + const size_t n_rows = layout.n_rows; + for (const auto& layer : layers) { + const size_t rc = layer.row_count; + if (rc == 0) { continue; } + const size_t batch = std::max(1, (size_t{64} << 20) / (sizeof(uint32_t) + urow)); + std::vector nodes(std::min(batch, rc)); + std::vector links(std::min(batch, rc) * urow); + for (size_t s = 0; s < rc; s += batch) { + const size_t cur = std::min(batch, rc - s); + cuvs::util::read_large_file( + artifact_fd, + nodes.data(), + cur * sizeof(uint32_t), + layout.upper_nodes_offset + (layer.node_offset + s) * sizeof(uint32_t)); + cuvs::util::read_large_file(artifact_fd, + links.data(), + cur * urow, + layout.upper_links_offset + (layer.link_offset + s) * urow); + for (size_t r = 0; r < cur; ++r) { + const size_t id = static_cast(nodes[r]); + RAFT_EXPECTS(id < n_rows, "Invalid upper-layer node id in layered HNSW artifact"); + RAFT_EXPECTS(layer.level <= static_cast(levels_u8[id]), + "Layered HNSW artifact references a node at an invalid upper level"); + cb(id, layer.level, links.data() + r * urow); + } + } + } +} + +// Emits the level-0 region `[link block | vector | label]` per ID, in increasing ID order, reading +// the dataset sequentially and writing the output sequentially. `row_ptr(id)` returns a pointer to +// the ID's level-0 link block in a caller-owned, bounded buffer. +template +void emit_hnswlib_base_records(const npy_file& dataset_file, + const cuvs::util::file_descriptor& output_fd, + const hnswlib_materialize_layout& layout, + int num_threads, + size_t id_begin, + size_t id_end, + std::vector& out_buffer, + DatasetBatch& dataset_batch, + RowPtr&& row_ptr) +{ + const size_t spe = layout.size_data_per_element; + const size_t row = layout.base_link_row_bytes; + const size_t data_size = layout.data_size; + const size_t out_batch_rows = static_cast(dataset_batch.extent(0)); + for (size_t s = id_begin; s < id_end; s += out_batch_rows) { + const size_t cur = std::min(out_batch_rows, id_end - s); + cuvs::util::read_large_file(dataset_file.fd, + dataset_batch.data_handle(), + cur * data_size, + dataset_file.header_size + s * data_size); +#pragma omp parallel for num_threads(num_threads) + for (int64_t k = 0; k < static_cast(cur); ++k) { + const size_t id = s + static_cast(k); + char* rec = out_buffer.data() + static_cast(k) * spe; + std::memcpy(rec, row_ptr(id), row); + std::memcpy(rec + layout.offset_data, + dataset_batch.data_handle() + static_cast(k) * layout.dim, + data_size); + const hnswlib::labeltype label = static_cast(id); + std::memcpy(rec + layout.label_offset, &label, sizeof(label)); + } + cuvs::util::write_large_file( + output_fd, out_buffer.data(), cur * spe, layout.base_region_offset + s * spe); + } +} + +// Phase 1 + 2: reorder the base topology to original-ID order and emit the level-0 region. +template +void materialize_hnswlib_base_region(const cuvs::util::file_descriptor& artifact_fd, + const npy_file& dataset_file, + const cuvs::util::file_descriptor& output_fd, + const layered_hnsw_file_metadata& metadata, + const hnswlib_materialize_layout& layout, + size_t budget_bytes, + int num_threads, + const std::filesystem::path& tmp_dir) +{ + const size_t n_rows = layout.n_rows; + const size_t row = layout.base_link_row_bytes; + + const size_t out_batch_rows = + std::max(1, (size_t{64} << 20) / layout.size_data_per_element); + auto dataset_batch = raft::make_host_matrix(static_cast(out_batch_rows), + static_cast(metadata.dim)); + std::vector out_buffer(out_batch_rows * layout.size_data_per_element); + + // Tracks that every original ID in [0, n_rows) is produced exactly once by the reorder. The + // section-size invariant (validated in Phase 0) guarantees exactly n_rows base-node records, so + // a malformed artifact with a missing/duplicate ID would otherwise leave a zero-initialized link + // row in the output and silently corrupt the graph; catch it instead of shipping a bad index. + std::vector seen(n_rows, 0); + auto verify_full_coverage = [&]() { + const auto covered = static_cast(std::count(seen.begin(), seen.end(), uint8_t{1})); + RAFT_EXPECTS(covered == n_rows, + "Layered HNSW base nodes cover %zu of %zu original ids (missing or duplicate ids)", + covered, + n_rows); + }; + + const auto base_start_time = std::chrono::steady_clock::now(); + // Single in-memory pass when the whole base topology fits the budget (no temporary files). + if (budget_bytes >= metadata.base_links_bytes) { + std::vector ordered_base(metadata.base_links_bytes); + scatter_layered_base_links(artifact_fd, layout, 1, [&](size_t id, const char* link_row) { + std::memcpy(ordered_base.data() + id * row, link_row, row); + seen[id] = 1; + }); + verify_full_coverage(); + emit_hnswlib_base_records(dataset_file, + output_fd, + layout, + num_threads, + 0, + n_rows, + out_buffer, + dataset_batch, + [&](size_t id) { return ordered_base.data() + id * row; }); + RAFT_LOG_INFO("hnswlib materialize: base region written (single-pass) in %ld ms (%.2f GiB)", + elapsed_ms_since(base_start_time), + to_gib(layout.size_data_per_element * n_rows)); + return; + } + + // Bucketed reorder through temporary files for a hard memory budget. + const size_t record_bytes = sizeof(IdxT) + row; + size_t rows_per_bucket = std::max(1, budget_bytes / record_bytes); + rows_per_bucket = std::min(rows_per_bucket, n_rows); + const size_t num_buckets = (n_rows + rows_per_bucket - 1) / rows_per_bucket; + RAFT_LOG_INFO( + "hnswlib materialize: base region uses %zu buckets (rows/bucket=%zu, budget=%.2f GiB)", + num_buckets, + rows_per_bucket, + to_gib(budget_bytes)); + + id_record_spiller spiller( + tmp_dir / "base", num_buckets, rows_per_bucket, record_bytes, budget_bytes / 2); + // The spiller's per-bucket mutexes make add() safe under the parallel scatter; each thread stages + // the record in its own (thread_local) buffer before the routed append. + scatter_layered_base_links(artifact_fd, layout, 1, [&](size_t id, const char* link_row) { + thread_local std::vector rec; + rec.resize(record_bytes); + const IdxT id32 = static_cast(id); + std::memcpy(rec.data(), &id32, sizeof(IdxT)); + std::memcpy(rec.data() + sizeof(IdxT), link_row, row); + spiller.add(id / rows_per_bucket, rec.data()); + seen[id] = 1; + }); + spiller.finish_writes(); + verify_full_coverage(); + + std::vector bucket_rows; + std::vector chunk; + for (size_t b = 0; b < num_buckets; ++b) { + const size_t id_begin = b * rows_per_bucket; + const size_t id_end = std::min(n_rows, id_begin + rows_per_bucket); + const size_t rib = id_end - id_begin; + bucket_rows.assign(rib * row, 0); + spiller.replay(b, chunk, [&](const char* r) { + IdxT id32; + std::memcpy(&id32, r, sizeof(IdxT)); + const size_t id = static_cast(id32); + std::memcpy(bucket_rows.data() + (id - id_begin) * row, r + sizeof(IdxT), row); + }); + emit_hnswlib_base_records( + dataset_file, + output_fd, + layout, + num_threads, + id_begin, + id_end, + out_buffer, + dataset_batch, + [&](size_t id) { return bucket_rows.data() + (id - id_begin) * row; }); + } + RAFT_LOG_INFO("hnswlib materialize: base region written (%zu buckets) in %ld ms (%.2f GiB)", + num_buckets, + elapsed_ms_since(base_start_time), + to_gib(layout.size_data_per_element * n_rows)); +} + +// Phase 3: transpose the upper layers into per-element link lists and emit the upper region. +inline void materialize_hnswlib_upper_region(const cuvs::util::file_descriptor& artifact_fd, + const cuvs::util::file_descriptor& output_fd, + const layered_hnsw_file_metadata& metadata, + const hnswlib_materialize_layout& layout, + const std::vector& levels_u8, + size_t budget_bytes, + const std::filesystem::path& tmp_dir) +{ + const size_t n_rows = layout.n_rows; + const size_t urow = layout.upper_link_row_bytes; + const auto upper_start_time = std::chrono::steady_clock::now(); + + // Sequential writer over the (variable-length) upper region. + size_t write_off = layout.upper_region_offset; + std::vector out_buffer; + out_buffer.reserve((size_t{64} << 20) + urow + sizeof(int)); + auto flush_out = [&]() { + if (out_buffer.empty()) { return; } + cuvs::util::write_large_file(output_fd, out_buffer.data(), out_buffer.size(), write_off); + write_off += out_buffer.size(); + out_buffer.clear(); + }; + auto append_element = [&](size_t level, const char* rows) { + const int link_list_size = level > 0 ? static_cast(level * urow) : 0; + const auto* p = reinterpret_cast(&link_list_size); + out_buffer.insert(out_buffer.end(), p, p + sizeof(int)); + if (level > 0) { out_buffer.insert(out_buffer.end(), rows, rows + level * urow); } + if (out_buffer.size() >= (size_t{64} << 20)) { flush_out(); } + }; + + const size_t fits_budget = metadata.upper_links_bytes + (n_rows + 1) * sizeof(size_t); + if (metadata.upper_links_bytes == 0 || fits_budget <= budget_bytes) { + // Pack all upper rows in original-ID order, then stream them out. + std::vector packed_start(n_rows + 1, 0); + for (size_t i = 0; i < n_rows; ++i) { + packed_start[i + 1] = packed_start[i] + static_cast(levels_u8[i]); + } + RAFT_EXPECTS(packed_start[n_rows] == metadata.upper_nodes_count, + "Layered HNSW upper rows (%zu) do not match upper_nodes_count (%zu)", + packed_start[n_rows], + metadata.upper_nodes_count); + std::vector packed(metadata.upper_links_bytes); + scatter_layered_upper_links(artifact_fd, + layout, + metadata.layers, + levels_u8, + [&](size_t id, size_t level, const char* link_row) { + const size_t dst_row = packed_start[id] + (level - 1); + std::memcpy(packed.data() + dst_row * urow, link_row, urow); + }); + size_t cursor = 0; + for (size_t id = 0; id < n_rows; ++id) { + const size_t level = static_cast(levels_u8[id]); + append_element(level, level > 0 ? packed.data() + cursor * urow : nullptr); + cursor += level; + } + flush_out(); + RAFT_LOG_INFO("hnswlib materialize: upper region written (single-pass) in %ld ms (%.2f GiB)", + elapsed_ms_since(upper_start_time), + to_gib(metadata.upper_links_bytes)); + return; + } + + // Bucketed upper transpose for a hard memory budget. + const size_t budget_half = std::max(1, budget_bytes / 2); + size_t num_buckets = + std::max(1, (metadata.upper_links_bytes + budget_half - 1) / budget_half); + size_t rows_per_bucket = std::max(1, (n_rows + num_buckets - 1) / num_buckets); + num_buckets = (n_rows + rows_per_bucket - 1) / rows_per_bucket; + RAFT_LOG_INFO("hnswlib materialize: upper region uses %zu buckets (rows/bucket=%zu)", + num_buckets, + rows_per_bucket); + + const size_t record_bytes = 2 * sizeof(uint32_t) + urow; // [id][level][row] + id_record_spiller spiller( + tmp_dir / "upper", num_buckets, rows_per_bucket, record_bytes, budget_half); + std::vector rec(record_bytes); + scatter_layered_upper_links(artifact_fd, + layout, + metadata.layers, + levels_u8, + [&](size_t id, size_t level, const char* link_row) { + const uint32_t id32 = static_cast(id); + const uint32_t lv32 = static_cast(level); + std::memcpy(rec.data(), &id32, sizeof(uint32_t)); + std::memcpy(rec.data() + sizeof(uint32_t), &lv32, sizeof(uint32_t)); + std::memcpy(rec.data() + 2 * sizeof(uint32_t), link_row, urow); + spiller.add(id / rows_per_bucket, rec.data()); + }); + spiller.finish_writes(); + + std::vector chunk; + for (size_t b = 0; b < num_buckets; ++b) { + const size_t id_begin = b * rows_per_bucket; + const size_t id_end = std::min(n_rows, id_begin + rows_per_bucket); + const size_t rib = id_end - id_begin; + std::vector local_start(rib + 1, 0); + for (size_t i = 0; i < rib; ++i) { + local_start[i + 1] = local_start[i] + static_cast(levels_u8[id_begin + i]); + } + std::vector packed(local_start[rib] * urow); + spiller.replay(b, chunk, [&](const char* r) { + uint32_t id32 = 0; + uint32_t lv32 = 0; + std::memcpy(&id32, r, sizeof(uint32_t)); + std::memcpy(&lv32, r + sizeof(uint32_t), sizeof(uint32_t)); + const size_t local = static_cast(id32) - id_begin; + const size_t dst_row = local_start[local] + (static_cast(lv32) - 1); + std::memcpy(packed.data() + dst_row * urow, r + 2 * sizeof(uint32_t), urow); + }); + size_t cursor = 0; + for (size_t i = 0; i < rib; ++i) { + const size_t level = static_cast(levels_u8[id_begin + i]); + append_element(level, level > 0 ? packed.data() + cursor * urow : nullptr); + cursor += level; + } + } + flush_out(); + RAFT_LOG_INFO("hnswlib materialize: upper region written (%zu buckets) in %ld ms (%.2f GiB)", + num_buckets, + elapsed_ms_since(upper_start_time), + to_gib(metadata.upper_links_bytes)); +} + +// Materialize a layered HNSW artifact + dataset into a standard hnswlib index file on disk, using +// bounded host memory and sequential disk I/O. +template +void materialize_layered_to_hnswlib_on_disk(raft::resources const& res, + const cuvs::util::file_descriptor& artifact_fd, + const cuvs::neighbors::hnsw::materialize_params& params, + const std::string& layered_artifact_path, + const std::string& output_path, + int dim, + cuvs::distance::DistanceType metric) +{ + static_assert(std::is_same_v, "Layered HNSW artifacts store ids as uint32_t"); + common::nvtx::range fun_scope("hnsw::materialize_to_hnswlib"); + const auto total_start_time = std::chrono::steady_clock::now(); + + // ---- Phase 0: read and validate the artifact header, descriptors and levels ---- + // Shared with the in-memory deserialize path so the format/offset arithmetic and every + // size-invariant check live in one place; this also bounds the buffer allocations below. + auto view = + read_and_validate_layered_artifact(artifact_fd, layered_artifact_path, dim, metric); + auto& metadata = view.metadata; + auto& levels_u8 = view.levels; + const auto base_nodes_offset = view.base_nodes_offset; + const auto base_links_offset = view.base_links_offset; + const auto upper_nodes_offset = view.upper_nodes_offset; + const auto upper_links_offset = view.upper_links_offset; + RAFT_EXPECTS(!params.dataset_path.empty(), + "Layered HNSW materialization requires materialize_params.dataset_path"); + + const size_t n_rows = metadata.n_rows; + + // Validate the dataset header. + auto dataset_file = open_layered_dataset_file(params.dataset_path); + validate_npy_file(dataset_file, params.dataset_path, "Layered HNSW dataset"); + RAFT_EXPECTS(dataset_file.shape.size() == 2 && dataset_file.shape[0] == n_rows && + dataset_file.shape[1] == metadata.dim, + "Layered HNSW dataset shape mismatch: artifact rows=%zu dim=%zu, dataset path=%s", + n_rows, + metadata.dim, + params.dataset_path.c_str()); + + // Retrieve the hnswlib layout constants from a dummy (single-element) index. + auto hnsw_index = std::make_unique>(dim, metric, HnswHierarchy::CPU); + auto appr_algo = std::make_unique::type>>( + hnsw_index->get_space(), 1, metadata.M, metadata.ef_construction); + const size_t data_size = static_cast(dim) * sizeof(T); + 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"); + RAFT_EXPECTS(appr_algo->data_size_ == data_size, "Layered HNSW data size mismatch"); + RAFT_EXPECTS(appr_algo->offsetData_ == appr_algo->size_links_level0_, + "Unexpected hnswlib data offset"); + RAFT_EXPECTS(appr_algo->label_offset_ == appr_algo->size_links_level0_ + data_size, + "Unexpected hnswlib label offset"); + + // ---- Output layout ---- + const auto native_header = make_hnswlib_native_header(appr_algo->offsetLevel0_, + n_rows, + appr_algo->size_data_per_element_, + appr_algo->label_offset_, + appr_algo->offsetData_, + metadata.maxlevel, + metadata.enterpoint_node, + metadata.maxM, + metadata.maxM0, + metadata.M, + metadata.mult, + metadata.ef_construction); + hnswlib_materialize_layout layout; + layout.n_rows = n_rows; + layout.dim = metadata.dim; + layout.base_link_row_bytes = metadata.base_link_row_bytes; + layout.upper_link_row_bytes = metadata.upper_link_row_bytes; + layout.size_data_per_element = appr_algo->size_data_per_element_; + layout.offset_data = appr_algo->offsetData_; + layout.label_offset = appr_algo->label_offset_; + layout.data_size = data_size; + layout.base_region_offset = native_header.size(); + layout.base_nodes_offset = base_nodes_offset; + layout.base_links_offset = base_links_offset; + layout.upper_nodes_offset = upper_nodes_offset; + layout.upper_links_offset = upper_links_offset; + + const size_t base_region_bytes = n_rows * layout.size_data_per_element; + layout.upper_region_offset = layout.base_region_offset + base_region_bytes; + + size_t upper_region_bytes = 0; + for (size_t i = 0; i < n_rows; ++i) { + upper_region_bytes += sizeof(int); + const size_t level = static_cast(levels_u8[i]); + if (level > 0) { upper_region_bytes += level * layout.upper_link_row_bytes; } + } + const size_t final_size = layout.upper_region_offset + upper_region_bytes; + + exclusive_hnsw_temp_file output(output_path); + cuvs::util::file_descriptor output_fd(output.temporary_path(), O_RDWR); + cuvs::util::preallocate_file(output_fd, final_size); + cuvs::util::write_large_file(output_fd, native_header.data(), native_header.size(), 0); + RAFT_LOG_INFO( + "hnswlib materialize: writing index (rows=%zu dim=%zu, output=%.2f GiB, dataset=%s)", + n_rows, + metadata.dim, + to_gib(final_size), + params.dataset_path.c_str()); + + auto num_threads = + params.num_threads == 0 ? cuvs::core::omp::get_max_threads() : params.num_threads; + // max_host_memory_gb <= 0 => no host-memory cap: both regions take their single in-memory pass + // and create no temporary files. Otherwise reserve headroom for the fixed-size streaming I/O + // buffers (dataset batch, output buffer, per-call scatter read buffers, each ~64 MiB) so peak + // host memory stays close to the requested budget rather than overshooting it by the buffer + // overhead. + size_t budget_bytes = std::numeric_limits::max(); + if (params.max_host_memory_gb > 0.0) { + const size_t requested = + std::max(1, static_cast(params.max_host_memory_gb * (size_t{1} << 30))); + constexpr size_t kStreamingBufferReserve = size_t{4} * (size_t{64} << 20); // ~256 MiB + budget_bytes = requested > kStreamingBufferReserve ? requested - kStreamingBufferReserve + : std::max(1, requested / 2); + } + + std::filesystem::path tmp_dir = output.temporary_path(); + tmp_dir += ".materialize_tmp"; + scoped_directory_cleanup workspace_cleanup(tmp_dir); + + materialize_hnswlib_base_region( + artifact_fd, dataset_file, output_fd, metadata, layout, budget_bytes, num_threads, tmp_dir); + materialize_hnswlib_upper_region( + artifact_fd, output_fd, metadata, layout, levels_u8, budget_bytes, tmp_dir); + + output_fd.close(); + output.publish_replace(); + + RAFT_LOG_INFO("hnswlib materialize: completed in %ld ms (output %.2f GiB)", + elapsed_ms_since(total_start_time), + to_gib(final_size)); +} + +// Dispatch using the external dataset dtype. The layered artifact records the graph-construction +// dtype, which may differ when a quantized graph is materialized with the original vectors. +inline void materialize_layered_to_hnswlib_on_disk_dispatch( + raft::resources const& res, + const cuvs::neighbors::hnsw::materialize_params& params, + const std::string& layered_artifact_path, + const std::string& output_path, + int dim, + cuvs::distance::DistanceType metric) +{ + RAFT_EXPECTS(!params.dataset_path.empty(), + "Layered HNSW materialization requires materialize_params.dataset_path"); + cuvs::util::file_descriptor artifact_fd(layered_artifact_path, O_RDONLY); + const auto dtype = layered_dataset_dtype(params.dataset_path); + switch (dtype) { + case layered_hnsw_dtype::float32: + materialize_layered_to_hnswlib_on_disk( + res, artifact_fd, params, layered_artifact_path, output_path, dim, metric); + break; + case layered_hnsw_dtype::float16: + materialize_layered_to_hnswlib_on_disk( + res, artifact_fd, params, layered_artifact_path, output_path, dim, metric); + break; + case layered_hnsw_dtype::uint8: + materialize_layered_to_hnswlib_on_disk( + res, artifact_fd, params, layered_artifact_path, output_path, dim, metric); + break; + case layered_hnsw_dtype::int8: + materialize_layered_to_hnswlib_on_disk( + res, artifact_fd, params, layered_artifact_path, output_path, dim, metric); + break; + default: + RAFT_FAIL("Unsupported layered HNSW dataset dtype or filename: %s", + params.dataset_path.c_str()); + } +} + template void deserialize(raft::resources const& res, const index_params& params, diff --git a/cpp/src/neighbors/hnsw.cpp b/cpp/src/neighbors/hnsw.cpp index 026f8a9bdd..7a07d7fd34 100644 --- a/cpp/src/neighbors/hnsw.cpp +++ b/cpp/src/neighbors/hnsw.cpp @@ -173,4 +173,17 @@ CUVS_INST_HNSW_SERIALIZE(int8_t); #undef CUVS_INST_HNSW_SERIALIZE +// The element data type is inferred from the external dataset; the dispatcher selects the typed +// implementation, instantiating the float/half/uint8_t/int8_t materialize paths in this TU. +void materialize_to_hnswlib(raft::resources const& res, + const materialize_params& params, + const std::string& layered_artifact_path, + const std::string& output_path, + int dim, + cuvs::distance::DistanceType metric) +{ + detail::materialize_layered_to_hnswlib_on_disk_dispatch( + res, params, layered_artifact_path, output_path, dim, metric); +} + } // namespace cuvs::neighbors::hnsw diff --git a/cpp/src/util/file_io.cpp b/cpp/src/util/file_io.cpp index 0b777a53cb..b524ad9341 100644 --- a/cpp/src/util/file_io.cpp +++ b/cpp/src/util/file_io.cpp @@ -13,10 +13,15 @@ #include #include #include +#include #include #include #include +#include +#include +#include + namespace cuvs::util { namespace { @@ -148,6 +153,26 @@ void write_large_file_posix(const file_descriptor& fd, } // namespace +void preallocate_file(const file_descriptor& fd, const size_t total_bytes) +{ + if (total_bytes == 0) { return; } + RAFT_EXPECTS(fd.is_valid(), "File descriptor must be valid"); + RAFT_EXPECTS(total_bytes <= static_cast(std::numeric_limits::max()), + "Requested file size exceeds the POSIX offset range"); + const int rc = posix_fallocate(fd.get(), 0, static_cast(total_bytes)); + if (rc == 0) { return; } + // Some filesystems (tmpfs, certain NFS/overlay mounts) do not support preallocation; fall back + // to ftruncate so a valid output location is still usable. + if (rc == EOPNOTSUPP || rc == EINVAL || rc == ENOSYS) { + RAFT_EXPECTS(ftruncate(fd.get(), static_cast(total_bytes)) == 0, + "Failed to pre-size file %s via ftruncate: %s", + fd.get_path().c_str(), + strerror(errno)); + return; + } + RAFT_FAIL("Failed to pre-allocate file %s: %s", fd.get_path().c_str(), strerror(rc)); +} + void read_large_file(const file_descriptor& fd, void* dest_ptr, const size_t total_bytes, diff --git a/cpp/tests/neighbors/ann_hnsw_ace.cuh b/cpp/tests/neighbors/ann_hnsw_ace.cuh index 90a3e46caf..f68c45c9e9 100644 --- a/cpp/tests/neighbors/ann_hnsw_ace.cuh +++ b/cpp/tests/neighbors/ann_hnsw_ace.cuh @@ -24,10 +24,12 @@ #include #include #include +#include #include #include #include #include +#include #include @@ -939,6 +941,246 @@ class AnnHnswAceTest : public ::testing::TestWithParam { } } + void testHnswAceLayeredMaterializeToHnswlib() + { + 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_materialize_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 hnsw_index = + hnsw::build(handle_, hnsw_params, raft::make_const_mdspan(database_host.view())); + ASSERT_NE(hnsw_index, nullptr); + const auto artifact_path = hnsw_index->file_path(); + ASSERT_FALSE(artifact_path.empty()); + + hnsw::search_params search_params; + search_params.ef = std::max(ps.ef_construction, ps.k * 2); + search_params.num_threads = 1; + + 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); + + // Reference: load the layered artifact in RAM and search. + std::vector indexes_layered(queries_size); + std::vector distances_layered(queries_size); + { + hnsw::index* layered_index = nullptr; + hnsw::deserialize(handle_, artifact_path, dataset_file, &layered_index); + ASSERT_NE(layered_index, nullptr); + std::unique_ptr> layered_guard(layered_index); + hnsw::search(handle_, + search_params, + *layered_guard, + queries_host.view(), + indexes_hnsw_host.view(), + distances_hnsw_host.view()); + for (size_t i = 0; i < queries_size; i++) { + indexes_layered[i] = static_cast(indexes_hnsw_host.data_handle()[i]); + distances_layered[i] = distances_hnsw_host.data_handle()[i]; + } + } + + // Materialize to a standard hnswlib index file twice: single in-memory pass and bucketed. + const auto out_single = (std::filesystem::path(temp_dir) / "materialized_single.bin").string(); + const auto out_bucketed = + (std::filesystem::path(temp_dir) / "materialized_bucketed.bin").string(); + + { + std::ofstream stale_output(out_single, std::ios::binary); + stale_output << "stale"; + } + + hnsw::materialize_params materialize_single; + materialize_single.dataset_path = dataset_file; + materialize_single.max_host_memory_gb = 0; // single in-memory reorder pass + hnsw::materialize_to_hnswlib( + handle_, materialize_single, artifact_path, out_single, ps.dim, ps.metric); + + hnsw::materialize_params materialize_bucketed; + materialize_bucketed.dataset_path = dataset_file; + materialize_bucketed.max_host_memory_gb = 0.00003; // tiny budget forces bucketed base + upper + hnsw::materialize_to_hnswlib( + handle_, materialize_bucketed, artifact_path, out_bucketed, ps.dim, ps.metric); + + ASSERT_TRUE(std::filesystem::is_regular_file(out_single)); + ASSERT_TRUE(std::filesystem::is_regular_file(out_bucketed)); + + // Determinism: single-pass and bucketed outputs must be byte-identical. + { + const auto read_all = [](const std::string& path) { + std::ifstream stream(path, std::ios::binary); + return std::vector((std::istreambuf_iterator(stream)), + std::istreambuf_iterator()); + }; + const auto bytes_single = read_all(out_single); + const auto bytes_bucketed = read_all(out_bucketed); + ASSERT_FALSE(bytes_single.empty()); + EXPECT_EQ(bytes_single.size(), bytes_bucketed.size()); + EXPECT_TRUE(bytes_single == bytes_bucketed) + << "Single-pass and bucketed materialized outputs differ"; + } + + if constexpr (std::is_same_v) { + auto float_database = raft::make_host_matrix(ps.n_rows, ps.dim); + std::transform(database_host.data_handle(), + database_host.data_handle() + ps.n_rows * ps.dim, + float_database.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, + float_database.data_handle(), + static_cast(ps.n_rows) * ps.dim * sizeof(float), + float_dataset_header_size); + + hnsw::materialize_params mixed_params; + mixed_params.dataset_path = float_dataset_file; + const auto mixed_output = + (std::filesystem::path(temp_dir) / "materialized_float.bin").string(); + hnsw::materialize_to_hnswlib( + handle_, mixed_params, artifact_path, mixed_output, ps.dim, ps.metric); + + hnsw::index_params mixed_cpu_params; + mixed_cpu_params.metric = ps.metric; + mixed_cpu_params.hierarchy = hnsw::HnswHierarchy::CPU; + hnsw::index* mixed_cpu_index = nullptr; + hnsw::deserialize( + handle_, mixed_cpu_params, mixed_output, ps.dim, ps.metric, &mixed_cpu_index); + ASSERT_NE(mixed_cpu_index, nullptr); + std::unique_ptr> mixed_cpu_guard(mixed_cpu_index); + } + + // Load the materialized file as a standard (CPU) hnswlib index and search. + hnsw::index_params cpu_params; + cpu_params.metric = ps.metric; + cpu_params.hierarchy = hnsw::HnswHierarchy::CPU; + hnsw::index* cpu_index = nullptr; + hnsw::deserialize(handle_, cpu_params, out_single, ps.dim, ps.metric, &cpu_index); + ASSERT_NE(cpu_index, nullptr); + std::unique_ptr> cpu_guard(cpu_index); + + hnsw::search(handle_, + search_params, + *cpu_guard, + queries_host.view(), + indexes_hnsw_host.view(), + distances_hnsw_host.view()); + + std::vector indexes_cpu(queries_size); + std::vector distances_cpu(queries_size); + for (size_t i = 0; i < queries_size; i++) { + indexes_cpu[i] = static_cast(indexes_hnsw_host.data_handle()[i]); + distances_cpu[i] = distances_hnsw_host.data_handle()[i]; + } + + EXPECT_TRUE(cuvs::neighbors::eval_neighbours(indexes_naive, + indexes_cpu, + distances_naive, + distances_cpu, + ps.n_queries, + ps.k, + 0.003, + ps.min_recall)) + << "Materialized hnswlib index failed recall check vs. ground truth"; + + // The materialized CPU index represents the same graph and vectors as the layered artifact, + // so its search results must match the in-memory layered path. + EXPECT_TRUE(cuvs::neighbors::eval_neighbours(indexes_layered, + indexes_cpu, + distances_layered, + distances_cpu, + ps.n_queries, + ps.k, + 0.003, + 0.99)) + << "Materialized hnswlib index disagrees with the in-memory layered path"; + + hnsw::materialize_params bad_params; + bad_params.dataset_path = dataset_file; + const auto err_out = (std::filesystem::path(temp_dir) / "err.bin").string(); + EXPECT_THROW(hnsw::materialize_to_hnswlib( + handle_, bad_params, artifact_path, err_out, ps.dim + 1, ps.metric), + std::exception); + EXPECT_FALSE(std::filesystem::exists(err_out)); + + const auto wrong_metric = ps.metric == cuvs::distance::DistanceType::L2Expanded + ? cuvs::distance::DistanceType::InnerProduct + : cuvs::distance::DistanceType::L2Expanded; + EXPECT_THROW(hnsw::materialize_to_hnswlib( + handle_, bad_params, artifact_path, err_out, ps.dim, wrong_metric), + std::exception); + EXPECT_FALSE(std::filesystem::exists(err_out)); + + hnsw::materialize_params missing_dataset; + missing_dataset.dataset_path.clear(); + EXPECT_THROW(hnsw::materialize_to_hnswlib( + handle_, missing_dataset, artifact_path, err_out, ps.dim, ps.metric), + std::exception); + EXPECT_FALSE(std::filesystem::exists(err_out)); + } + void SetUp() override { database_dev.resize(((size_t)ps.n_rows) * ps.dim, stream_); 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 6cadb07646..6331f5e33e 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 @@ -74,4 +74,14 @@ INSTANTIATE_TEST_CASE_P(AnnHnswInmemSpillTest, AnnHnswInmemSpillTest_float, ::testing::ValuesIn(hnsw_inmem_spill_inputs)); +typedef AnnHnswAceTest AnnHnswAceMaterializeTest_float; +TEST_P(AnnHnswAceMaterializeTest_float, AnnHnswAceLayeredMaterializeToHnswlib) +{ + this->testHnswAceLayeredMaterializeToHnswlib(); +} + +INSTANTIATE_TEST_CASE_P(AnnHnswAceMaterializeTest, + AnnHnswAceMaterializeTest_float, + ::testing::ValuesIn(hnsw_ace_layered_inputs)); + } // namespace cuvs::neighbors::hnsw 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 67d72c81d8..c5b64607ce 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 @@ -64,4 +64,14 @@ INSTANTIATE_TEST_CASE_P(AnnHnswInmemSpillTest, AnnHnswInmemSpillTest_half, ::testing::ValuesIn(hnsw_inmem_spill_inputs)); +typedef AnnHnswAceTest AnnHnswAceMaterializeTest_half; +TEST_P(AnnHnswAceMaterializeTest_half, AnnHnswAceLayeredMaterializeToHnswlib) +{ + this->testHnswAceLayeredMaterializeToHnswlib(); +} + +INSTANTIATE_TEST_CASE_P(AnnHnswAceMaterializeTest, + AnnHnswAceMaterializeTest_half, + ::testing::ValuesIn(hnsw_ace_layered_inputs)); + } // namespace cuvs::neighbors::hnsw 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 52202d7a22..5ef0b64427 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 @@ -66,4 +66,14 @@ INSTANTIATE_TEST_CASE_P(AnnHnswInmemSpillTest, AnnHnswInmemSpillTest_int8_t, ::testing::ValuesIn(hnsw_inmem_spill_inputs)); +typedef AnnHnswAceTest AnnHnswAceMaterializeTest_int8_t; +TEST_P(AnnHnswAceMaterializeTest_int8_t, AnnHnswAceLayeredMaterializeToHnswlib) +{ + this->testHnswAceLayeredMaterializeToHnswlib(); +} + +INSTANTIATE_TEST_CASE_P(AnnHnswAceMaterializeTest, + AnnHnswAceMaterializeTest_int8_t, + ::testing::ValuesIn(hnsw_ace_layered_inputs)); + } // namespace cuvs::neighbors::hnsw 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 a0825c0c9d..db73bb47e5 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 @@ -66,4 +66,14 @@ INSTANTIATE_TEST_CASE_P(AnnHnswInmemSpillTest, AnnHnswInmemSpillTest_uint8_t, ::testing::ValuesIn(hnsw_inmem_spill_inputs)); +typedef AnnHnswAceTest AnnHnswAceMaterializeTest_uint8_t; +TEST_P(AnnHnswAceMaterializeTest_uint8_t, AnnHnswAceLayeredMaterializeToHnswlib) +{ + this->testHnswAceLayeredMaterializeToHnswlib(); +} + +INSTANTIATE_TEST_CASE_P(AnnHnswAceMaterializeTest, + AnnHnswAceMaterializeTest_uint8_t, + ::testing::ValuesIn(hnsw_ace_layered_inputs)); + } // namespace cuvs::neighbors::hnsw diff --git a/examples/cpp/src/hnsw_ace_layered_example.cu b/examples/cpp/src/hnsw_ace_layered_example.cu index 14de426605..92a4264c65 100644 --- a/examples/cpp/src/hnsw_ace_layered_example.cu +++ b/examples/cpp/src/hnsw_ace_layered_example.cu @@ -7,16 +7,19 @@ // 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. +// HnswOutputFormat::GRAPH_ONLY builds a GPU hierarchy as a topology-only artifact on disk. It emits +// one 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: +// This example demonstrates how to build a layered HNSW index with ACE and turn it into a standard +// hnswlib index for in-memory search: // // 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. +// 2. Build a single-file layered HNSW artifact with ACE using hnsw::build. +// 3. Materialize the layered artifact into a standard hnswlib index file on disk using +// hnsw::materialize_to_hnswlib (disk-to-disk, never holding the full index in host memory). +// 4. Read the materialized hnswlib index into memory using hnsw::deserialize (hierarchy = CPU). +// 5. Search the in-memory HNSW index. // // Layered-on-disk layout: // @@ -26,26 +29,27 @@ // 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. +// The transferred index artifact is topology-only. The dataset is loaded locally during +// materialization from hnsw::materialize_params::dataset_path. The loader supports .npy and ANN +// benchmark *.bin datasets; this example writes a local dataset .npy only to make the demo +// self-contained. The materialized hnswlib index file is self-contained (it embeds the vectors), +// so reading it back needs no dataset path. #include +#include +#include #include #include #include #include #include #include +#include #include #include #include +#include #include #include @@ -64,6 +68,18 @@ namespace { constexpr const char* kBuildDir = "/tmp/hnsw_ace_layered"; +// Reports the wall-clock time of a callable in milliseconds. +template +double time_ms(F&& fn) +{ + const auto start = std::chrono::steady_clock::now(); + fn(); + return std::chrono::duration(std::chrono::steady_clock::now() - start) + .count(); +} + +double to_gib(double bytes) { return bytes / (1024.0 * 1024.0 * 1024.0); } + template std::string write_local_dataset(raft::host_matrix_view dataset, const std::string& path) @@ -75,9 +91,15 @@ std::string write_local_dataset(raft::host_matrix_view dataset return path; } -auto quantize_dataset(raft::device_resources const& dev_resources, - raft::host_matrix_view dataset_float) - -> raft::host_matrix +template +struct quantized_pair { + raft::host_matrix dataset; + raft::host_matrix queries; +}; + +quantized_pair quantize_dataset(raft::device_resources const& dev_resources, + raft::host_matrix_view dataset_float, + raft::host_matrix_view queries_float) { std::cout << " quantize_dataset: training scalar quantizer (float -> int8)" << std::endl; cuvs::preprocessing::quantize::scalar::params qp; @@ -87,7 +109,13 @@ auto quantize_dataset(raft::device_resources const& dev_resources, 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 queries_i8 = + raft::make_host_matrix(queries_float.extent(0), queries_float.extent(1)); + cuvs::preprocessing::quantize::scalar::transform( + dev_resources, quantizer, queries_float, queries_i8.view()); + + return {std::move(dataset_i8), std::move(queries_i8)}; } auto make_hnsw_ace_params(const std::string& build_dir) -> cuvs::neighbors::hnsw::index_params @@ -117,26 +145,72 @@ auto hnsw_build(raft::device_resources const& dev_resources, { using namespace cuvs::neighbors; - auto hnsw_index = hnsw::build(dev_resources, hnsw_params, dataset); + std::unique_ptr> hnsw_index; + const auto build_ms = + time_ms([&]() { 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; + const auto artifact_bytes = static_cast(std::filesystem::file_size(artifact_path)); + std::cout << " hnsw_build: layered artifact written to " << artifact_path << "\n" + << " hnsw_build: build wall time " << build_ms << " ms, artifact " + << to_gib(artifact_bytes) << " GiB" << std::endl; return artifact_path; } +// Materialize the layered artifact into a standard hnswlib index file on disk and time the +// disk-to-disk materialization. Returns the path to the native hnswlib index file. template -auto hnsw_deserialize(raft::device_resources const& dev_resources, +auto hnsw_materialize(raft::device_resources const& dev_resources, + const cuvs::neighbors::hnsw::index_params& hnsw_params, const std::string& artifact_path, - const std::string& dataset_path) - -> std::unique_ptr> + const std::string& dataset_path, + int64_t dim, + const std::string& output_path) -> std::string +{ + using namespace cuvs::neighbors; + + hnsw::materialize_params materialize_params; + materialize_params.dataset_path = dataset_path; + materialize_params.max_host_memory_gb = 0; // 0 => single in-memory reorder pass + materialize_params.num_threads = 0; // 0 => max threads + + const auto materialize_ms = time_ms([&]() { + hnsw::materialize_to_hnswlib(dev_resources, + materialize_params, + artifact_path, + output_path, + static_cast(dim), + hnsw_params.metric); + }); + + const auto native_bytes = static_cast(std::filesystem::file_size(output_path)); + std::cout << " hnsw_materialize: native hnswlib index written to " << output_path << "\n" + << " hnsw_materialize: wall time " << materialize_ms << " ms, output " + << to_gib(native_bytes) << " GiB" << std::endl; + return output_path; +} + +// Read the materialized hnswlib index into memory for search. The materialized file is a standard +// hnswlib index, so it is loaded with hierarchy == CPU and needs no dataset path (the file already +// embeds the vectors). +template +auto hnsw_load_native(raft::device_resources const& dev_resources, + const std::string& native_index_path, + cuvs::distance::DistanceType metric, + int64_t dim) -> 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); + hnsw::index_params load_params; + load_params.hierarchy = hnsw::HnswHierarchy::CPU; + load_params.metric = metric; + + hnsw::index* loaded_index = nullptr; + hnsw::deserialize( + dev_resources, load_params, native_index_path, static_cast(dim), metric, &loaded_index); + return std::unique_ptr>(loaded_index); } template @@ -190,6 +264,9 @@ int main() { raft::device_resources dev_resources; + // Surface the per-phase build/materialize timing logs (RAFT_LOG_INFO). + raft::default_logger().set_level(rapids_logger::level_enum::info); + 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); @@ -227,33 +304,47 @@ int main() std::filesystem::create_directories(kBuildDir); #if HNSW_ACE_LAYERED_USE_QUANTIZATION - auto dataset_i8 = quantize_dataset(dev_resources, dataset_host_view); + auto q = quantize_dataset(dev_resources, dataset_host_view, queries_host_view); auto dataset_i8_view = raft::make_host_matrix_view( - dataset_i8.data_handle(), n_samples, n_dim); + q.dataset.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; + const std::string native_index_path = std::string{kBuildDir} + "/hnsw_native.bin"; + + std::cout << "[stage 2] Build layered HNSW index 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 3] Materialize layered HNSW -> native hnswlib index" << std::endl; + hnsw_materialize( + dev_resources, hnsw_params, artifact_path, dataset_path, n_dim, native_index_path); + + std::cout << "[stage 4] Read materialized hnswlib index into memory" << std::endl; + auto hnsw_index = + hnsw_load_native(dev_resources, native_index_path, hnsw_params.metric, n_dim); - std::cout << "[stage 4] Search float HNSW index" << std::endl; + std::cout << "[stage 5] Search 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); + const std::string native_index_path = std::string{kBuildDir} + "/hnsw_native.bin"; + 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 3] Materialize layered HNSW -> native hnswlib index" << std::endl; + hnsw_materialize( + dev_resources, hnsw_params, artifact_path, dataset_path, n_dim, native_index_path); + + std::cout << "[stage 4] Read materialized hnswlib index into memory" << std::endl; + auto hnsw_index = + hnsw_load_native(dev_resources, native_index_path, hnsw_params.metric, n_dim); - std::cout << "[stage 4] Search HNSW index" << std::endl; + std::cout << "[stage 5] Search HNSW index" << std::endl; hnsw_search(dev_resources, *hnsw_index, queries_host_view); #endif diff --git a/fern/docs.yml b/fern/docs.yml index 77ee0dae2a..1daa52cf02 100644 --- a/fern/docs.yml +++ b/fern/docs.yml @@ -538,6 +538,8 @@ navigation: path: "./pages/java_api/java-api-com-nvidia-cuvs-hnswindex.md" - page: "HnswIndexParams" path: "./pages/java_api/java-api-com-nvidia-cuvs-hnswindexparams.md" + - page: "HnswMaterializeParams" + path: "./pages/java_api/java-api-com-nvidia-cuvs-hnswmaterializeparams.md" - page: "HnswQuery" path: "./pages/java_api/java-api-com-nvidia-cuvs-hnswquery.md" - page: "HnswSearchParams" diff --git a/fern/pages/c_api/c-api-neighbors-hnsw.md b/fern/pages/c_api/c-api-neighbors-hnsw.md index 4d46dbfa0e..4dd58c5335 100644 --- a/fern/pages/c_api/c-api-neighbors-hnsw.md +++ b/fern/pages/c_api/c-api-neighbors-hnsw.md @@ -487,3 +487,95 @@ NOTE: When hierarchy is `NONE`, the loaded hnswlib index is immutable, and only **Returns** [`cuvsError_t`](/api-reference/c-api-core-c-api#cuvserror-t) + +## Materialize a layered HNSW artifact to an hnswlib index + + +### cuvsHnswMaterializeParams + +Parameters for materializing a layered HNSW artifact into an hnswlib index on disk. + +```c +struct cuvsHnswMaterializeParams { + const char* dataset_path; + double max_host_memory_gb; + int num_threads; +}; +``` + +**Fields** + +| Name | Type | Description | +| --- | --- | --- | +| `dataset_path` | `const char*` | Local dataset path holding the original-ID-ordered vectors used to build the artifact.

Supported formats match layered deserialization: `.npy` and ANN benchmark `*.bin` files with a `[uint32 rows, uint32 cols]` header (`.fbin`, `.f16bin`, `.u8bin`, `.i8bin`). | +| `max_host_memory_gb` | `double` | Upper bound on host memory (in GiB) used for the base-topology reorder buffer.

When `<= 0`, the whole base topology is reordered in a single in-memory pass (no temporary files). When set, the base topology is reordered through bucketed temporary files so that peak host memory stays close to this budget. | +| `num_threads` | `int` | Number of host threads to use. When `0`, the maximum number of threads is used. | + + +### cuvsHnswMaterializeParamsCreate + +Allocate HNSW materialize params, and populate with default values + +```c +cuvsError_t cuvsHnswMaterializeParamsCreate(cuvsHnswMaterializeParams_t* params); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `params` | in | [`cuvsHnswMaterializeParams_t*`](/api-reference/c-api-neighbors-hnsw#cuvshnswmaterializeparams) | cuvsHnswMaterializeParams_t to allocate | + +**Returns** + +[`cuvsError_t`](/api-reference/c-api-core-c-api#cuvserror-t) + + +### cuvsHnswMaterializeParamsDestroy + +De-allocate HNSW materialize params + +```c +cuvsError_t cuvsHnswMaterializeParamsDestroy(cuvsHnswMaterializeParams_t params); +``` + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `params` | in | [`cuvsHnswMaterializeParams_t`](/api-reference/c-api-neighbors-hnsw#cuvshnswmaterializeparams) | cuvsHnswMaterializeParams_t to de-allocate | + +**Returns** + +[`cuvsError_t`](/api-reference/c-api-core-c-api#cuvserror-t) + + +### cuvsHnswMaterializeToHnswlib + +Materialize a layered HNSW artifact into a standard hnswlib index file on disk. + +```c +cuvsError_t cuvsHnswMaterializeToHnswlib(cuvsResources_t res, +cuvsHnswMaterializeParams_t params, +const char* layered_artifact_path, +const char* output_path, +int dim, +cuvsDistanceType metric); +``` + +Materializes a `GRAPH_ONLY` artifact (graph topology only, stored in ACE order) plus a local dataset into a standard hnswlib index file, without ever holding the full materialized index in host memory. The resulting file is compatible with the original hnswlib library and can be read back through `cuvsHnswDeserialize` with `hierarchy == CPU`. The element data type (`float`, `half`, `uint8_t` or `int8_t`) is inferred from the external dataset. GRAPH_ONLY artifacts are currently produced through the C++ API. + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `res` | in | [`cuvsResources_t`](/api-reference/c-api-core-c-api#cuvsresources-t) | cuvsResources_t opaque C handle | +| `params` | in | [`cuvsHnswMaterializeParams_t`](/api-reference/c-api-neighbors-hnsw#cuvshnswmaterializeparams) | cuvsHnswMaterializeParams_t materialization parameters | +| `layered_artifact_path` | in | `const char*` | path to the layered HNSW artifact | +| `output_path` | in | `const char*` | path to the hnswlib index file to write | +| `dim` | in | `int` | the dimension of the vectors in the index | +| `metric` | in | [`cuvsDistanceType`](/api-reference/c-api-distance-distance#cuvsdistancetype) | the distance metric used to build the index | + +**Returns** + +[`cuvsError_t`](/api-reference/c-api-core-c-api#cuvserror-t) diff --git a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md index cdf32501d4..df187301c6 100644 --- a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md +++ b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md @@ -1102,3 +1102,61 @@ index** index); **Returns** `void` + +## Materialize a layered HNSW artifact into an hnswlib index + + +### neighbors::hnsw::materialize_params + +Parameters for materializing a layered HNSW artifact into an hnswlib index on disk. + +```cpp +struct materialize_params { + std::string dataset_path; + double max_host_memory_gb; + int num_threads; +}; +``` + +**Fields** + +| Name | Type | Description | +| --- | --- | --- | +| `dataset_path` | `std::string` | Local dataset path holding the original-ID-ordered vectors used to build the artifact.

Supported formats match layered deserialization: `.npy` and ANN benchmark `*.bin` files with a `[uint32 rows, uint32 cols]` header (`.fbin`, `.f16bin`, `.u8bin`, `.i8bin`). | +| `max_host_memory_gb` | `double` | Upper bound on host memory (in GiB) used for the base-topology reorder buffer.

When `<= 0`, the whole base topology is reordered in a single in-memory pass (no temporary files). When set, the base topology is reordered through bucketed temporary files so that peak host memory stays close to this budget, at the cost of writing and re-reading the (small) base-topology section once. | +| `num_threads` | `int` | Number of host threads to use. When `0`, the maximum number of threads is used. | + + +### neighbors::hnsw::materialize_to_hnswlib + +Materialize a layered HNSW artifact into a standard hnswlib index file on disk. + +```cpp +void materialize_to_hnswlib(raft::resources const& res, +const materialize_params& params, +const std::string& layered_artifact_path, +const std::string& output_path, +int dim, +cuvs::distance::DistanceType metric); +``` + +Materializes a `GRAPH_ONLY` artifact (graph topology only, stored in ACE order) plus a local dataset into a standard hnswlib index file, without ever holding the full materialized index in host memory. The materialization reorders the base topology from ACE order to original-id order and interleaves the vectors, emitting the output with sequential disk IO. The resulting file is compatible with the original hnswlib library (`loadIndex`) and can be read back through `cuvs::neighbors::hnsw::deserialize` with `hierarchy == HnswHierarchy::CPU`. + +The element data type (`float`, `half`, `uint8_t` or `int8_t`) is inferred from the external dataset, so materialization supports an original dataset dtype that differs from the graph's construction dtype. + +Usage example: + +**Parameters** + +| Name | Direction | Type | Description | +| --- | --- | --- | --- | +| `res` | in | `raft::resources const&` | raft resources | +| `params` | in | [`const materialize_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-materialize-params) | materialization parameters (dataset path, host-memory budget, threads) | +| `layered_artifact_path` | in | `const std::string&` | path to the layered HNSW artifact | +| `output_path` | in | `const std::string&` | path to the hnswlib index file to write | +| `dim` | in | `int` | dimensions of the training dataset | +| `metric` | in | [`cuvs::distance::DistanceType`](/api-reference/cpp-api-distance-distance#distance-distancetype) | distance metric. Supported metrics ("L2Expanded", "InnerProduct") | + +**Returns** + +`void` diff --git a/fern/pages/java_api/index.md b/fern/pages/java_api/index.md index 3f0293025d..f94a35fb89 100644 --- a/fern/pages/java_api/index.md +++ b/fern/pages/java_api/index.md @@ -40,6 +40,7 @@ For the Apache Lucene codecs built on this API, see the [Lucene API Documentatio - [HnswAceParams](/api-reference/java-api-com-nvidia-cuvs-hnswaceparams) - [HnswIndex](/api-reference/java-api-com-nvidia-cuvs-hnswindex) - [HnswIndexParams](/api-reference/java-api-com-nvidia-cuvs-hnswindexparams) +- [HnswMaterializeParams](/api-reference/java-api-com-nvidia-cuvs-hnswmaterializeparams) - [HnswQuery](/api-reference/java-api-com-nvidia-cuvs-hnswquery) - [HnswSearchParams](/api-reference/java-api-com-nvidia-cuvs-hnswsearchparams) - [MultiPartitionCagraSearch](/api-reference/java-api-com-nvidia-cuvs-multipartitioncagrasearch) diff --git a/fern/pages/java_api/java-api-com-nvidia-cuvs-hnswindex.md b/fern/pages/java_api/java-api-com-nvidia-cuvs-hnswindex.md index 338dfdb0b3..8bbb6beb1e 100644 --- a/fern/pages/java_api/java-api-com-nvidia-cuvs-hnswindex.md +++ b/fern/pages/java_api/java-api-com-nvidia-cuvs-hnswindex.md @@ -125,6 +125,42 @@ A new HNSW index ready for search _Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java:70`_ +### materializeToHnswlib + +```java +static void materializeToHnswlib( CuVSResources resources, HnswMaterializeParams materializeParams, String layeredArtifactPath, String outputPath, int dim, HnswIndexParams.CuvsDistanceType metric) throws Throwable +``` + +Materializes a layered HNSW artifact into a standard hnswlib index file on +disk. + +Materializes a `GRAPH_ONLY` artifact (graph topology only, +stored in ACE order) plus a local dataset into a standard hnswlib index file, +without ever holding the full materialized index in host memory. The +resulting file is compatible with the original hnswlib library and can be read +back with `hierarchy == CPU`. The element data type is inferred from the +external dataset. GRAPH_ONLY artifacts are currently produced through the C++ +API. + +**Parameters** + +| Name | Description | +| --- | --- | +| `resources` | The CuVS resources | +| `materializeParams` | Materialization parameters (dataset path, host-memory budget, threads) | +| `layeredArtifactPath` | Path to the layered HNSW artifact | +| `outputPath` | Path to the hnswlib index file to write | +| `dim` | The dimension of the vectors in the index | +| `metric` | The distance metric used to build the index | + +**Throws** + +| Type | Description | +| --- | --- | +| `Throwable` | if an error occurs during materialization | + +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java:99`_ + ### from ```java @@ -144,7 +180,7 @@ needed. an instance of this Builder -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java:90`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java:129`_ ### withIndexParams @@ -165,7 +201,7 @@ Builder. An instance of this Builder. -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java:99`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java:138`_ ### build @@ -179,6 +215,6 @@ Builds and returns an instance of CagraIndex. an instance of CagraIndex -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java:106`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java:145`_ _Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java:17`_ diff --git a/fern/pages/java_api/java-api-com-nvidia-cuvs-hnswmaterializeparams.md b/fern/pages/java_api/java-api-com-nvidia-cuvs-hnswmaterializeparams.md new file mode 100644 index 0000000000..746ab0ecf8 --- /dev/null +++ b/fern/pages/java_api/java-api-com-nvidia-cuvs-hnswmaterializeparams.md @@ -0,0 +1,159 @@ +--- +slug: api-reference/java-api-com-nvidia-cuvs-hnswmaterializeparams +--- + +# HnswMaterializeParams + +_Java package: `com.nvidia.cuvs`_ + +```java +public class HnswMaterializeParams +``` + +Parameters for materializing a layered HNSW artifact into a standard hnswlib +index file on disk. + +## Public Members + +### getDatasetPath + +```java +public String getDatasetPath() +``` + +Gets the local dataset path holding the original-ID-ordered vectors used to +build the artifact. + +**Returns** + +the dataset path, or null if not set + +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java:30`_ + +### getMaxHostMemoryGb + +```java +public double getMaxHostMemoryGb() +``` + +Gets the upper bound on host memory (in GiB) used for the base-topology +reorder buffer. When `<= 0`, the whole base topology is reordered in a +single in-memory pass. + +**Returns** + +the max host memory in GiB (0 means a single in-memory pass) + +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java:41`_ + +### getNumThreads + +```java +public int getNumThreads() +``` + +Gets the number of host threads to use. When 0, the maximum number of +threads is used. + +**Returns** + +the number of threads + +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java:51`_ + +### Builder + +```java +public Builder() +``` + +Constructs this Builder. + +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java:78`_ + +### withDatasetPath + +```java +public Builder withDatasetPath(String datasetPath) +``` + +Sets the local dataset path holding the original-ID-ordered vectors used to +build the artifact. Supported formats match layered deserialization: +`.npy` and ANN benchmark `*.bin` files with a +`[uint32 rows, uint32 cols]` header (`.fbin`, `.f16bin`, +`.u8bin`, `.i8bin`). + +**Parameters** + +| Name | Description | +| --- | --- | +| `datasetPath` | the local dataset path | + +**Returns** + +an instance of Builder + +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java:90`_ + +### withMaxHostMemoryGb + +```java +public Builder withMaxHostMemoryGb(double maxHostMemoryGb) +``` + +Sets the upper bound on host memory (in GiB) used for the base-topology +reorder buffer. + +When `<= 0` (default), the whole base topology is reordered in a single +in-memory pass (no temporary files). When set, the base topology is reordered +through bucketed temporary files so that peak host memory stays close to this +budget. + +**Parameters** + +| Name | Description | +| --- | --- | +| `maxHostMemoryGb` | the max host memory in GiB | + +**Returns** + +an instance of Builder + +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java:107`_ + +### withNumThreads + +```java +public Builder withNumThreads(int numThreads) +``` + +Sets the number of host threads to use. When 0 (default), the maximum number +of threads is used. + +**Parameters** + +| Name | Description | +| --- | --- | +| `numThreads` | the number of threads | + +**Returns** + +an instance of Builder + +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java:119`_ + +### build + +```java +public HnswMaterializeParams build() +``` + +Builds an instance of `HnswMaterializeParams`. + +**Returns** + +an instance of `HnswMaterializeParams` + +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java:129`_ + +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java:13`_ diff --git a/fern/pages/java_api/java-api-com-nvidia-cuvs-spi-cuvsprovider.md b/fern/pages/java_api/java-api-com-nvidia-cuvs-spi-cuvsprovider.md index 6230de4ee7..96f32f83b7 100644 --- a/fern/pages/java_api/java-api-com-nvidia-cuvs-spi-cuvsprovider.md +++ b/fern/pages/java_api/java-api-com-nvidia-cuvs-spi-cuvsprovider.md @@ -286,6 +286,33 @@ A new HNSW index ready for search _Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:165`_ +### hnswMaterializeToHnswlib + +```java +default void hnswMaterializeToHnswlib( CuVSResources resources, HnswMaterializeParams materializeParams, String layeredArtifactPath, String outputPath, int dim, HnswIndexParams.CuvsDistanceType metric) throws Throwable +``` + +Materializes a layered HNSW artifact into a standard hnswlib index file on disk. + +**Parameters** + +| Name | Description | +| --- | --- | +| `resources` | The CuVS resources | +| `materializeParams` | Materialization parameters (dataset path, host-memory budget, threads) | +| `layeredArtifactPath` | Path to the layered HNSW artifact | +| `outputPath` | Path to the hnswlib index file to write | +| `dim` | The dimension of the vectors in the index | +| `metric` | The distance metric used to build the index | + +**Throws** + +| Type | Description | +| --- | --- | +| `Throwable` | if an error occurs during materialization | + +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:179`_ + ### newTieredIndexBuilder ```java @@ -294,7 +321,7 @@ TieredIndex.Builder newTieredIndexBuilder(CuVSResources cuVSResources) throws Un Creates a new TieredIndex Builder. -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:169`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:191`_ ### isCagraPaddedDataset @@ -327,7 +354,7 @@ true when the rows are already padded the way CAGRA requires | --- | --- | | `UnsupportedOperationException` | if this provider cannot answer | -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:186`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:208`_ ### mergeCagraIndexes @@ -357,7 +384,7 @@ A new merged CAGRA index | --- | --- | | `Throwable` | if an error occurs during the merge operation | -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:202`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:224`_ ### newFilterBitsetHandle @@ -374,7 +401,7 @@ Per-partition bit offsets are recomputed inside cuVS from the index sizes. | --- | --- | | `combinedLongs` | packed bitset words for a single partition | -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:211`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:233`_ ### searchCagraMultiPartition @@ -400,7 +427,7 @@ Searches multiple CAGRA index partitions for the global top-k nearest neighbors | --- | --- | | `Throwable` | if an error occurs during the search | -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:224`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:246`_ ### gpuInfoProvider @@ -410,7 +437,7 @@ GPUInfoProvider gpuInfoProvider() Returns a `GPUInfoProvider` to query the system for GPU related information -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:233`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:255`_ ### enableRMMPooledMemory @@ -429,7 +456,7 @@ This operation has a global effect, and will affect all resources on the current | `initialPoolSizePercent` | The initial pool size, in percentage of the total GPU memory | | `maxPoolSizePercent` | The maximum pool size, in percentage of the total GPU memory | -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:247`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:269`_ ### enableRMMManagedPooledMemory @@ -448,7 +475,7 @@ This operation has a global effect, and will affect all resources on the current | `initialPoolSizePercent` | The initial pool size, in percentage of the total GPU memory | | `maxPoolSizePercent` | The maximum pool size, in percentage of the total GPU memory | -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:257`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:279`_ ### enableRMMAsyncMemory @@ -463,7 +490,7 @@ on deallocation. This is especially beneficial when multiple CAGRA searches run on separate CUDA streams, because internal workspace allocations no longer serialize kernel launches. This operation has a global effect and will affect all resources on the current device. -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:267`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:289`_ ### resetRMMPooledMemory @@ -473,7 +500,7 @@ void resetRMMPooledMemory() Disables pooled memory on the current device, reverting back to the default setting. -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:270`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:292`_ ### provider @@ -483,7 +510,7 @@ static CuVSProvider provider() Retrieves the system-wide provider. -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:273`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:295`_ ### cagraIndexParamsFromHnswParams @@ -513,7 +540,7 @@ may be shifted along the curve right or left. See the heuristics descriptions fo A new CAGRA index parameters object -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:293`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:315`_ ### cagraIndexParamsFromDataset @@ -537,6 +564,6 @@ Create CAGRA index parameters heuristically tuned for a dataset. A new CAGRA index parameters object -_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:311`_ +_Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:333`_ _Source: `java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java:18`_ diff --git a/fern/pages/python_api/python-api-neighbors-hnsw.md b/fern/pages/python_api/python-api-neighbors-hnsw.md index d4578b81ad..14c4fd451f 100644 --- a/fern/pages/python_api/python-api-neighbors-hnsw.md +++ b/fern/pages/python_api/python-api-neighbors-hnsw.md @@ -192,6 +192,55 @@ def __init__(self, *, num_threads=0) def num_threads(self) ``` +## MaterializeParams + +```python +cdef class MaterializeParams +``` + +Parameters for materializing a layered HNSW artifact into an hnswlib +index on disk. + +**Parameters** + +| Name | Type | Description | +| --- | --- | --- | +| `dataset_path` | `string, default = None (optional)` | Local dataset path holding the original-ID-ordered vectors used to build the artifact. Supported formats match layered deserialization: `.npy` and ANN benchmark `*.bin` files with a `[uint32 rows, uint32 cols]` header (`.fbin`, `.f16bin`, `.u8bin`, `.i8bin`). | +| `max_host_memory_gb` | `float, default = 0 (optional)` | Upper bound on host memory (in GiB) used for the base-topology reorder buffer. When <= 0, the whole base topology is reordered in a single in-memory pass (no temporary files). When set, the base topology is reordered through bucketed temporary files so that peak host memory stays close to this budget. | +| `num_threads` | `int, default = 0 (optional)` | Number of host threads to use. When 0, the maximum number of threads is used. | + +**Constructor** + +```python +def __init__(self, *, dataset_path=None, max_host_memory_gb=0, num_threads=0) +``` + +**Members** + +| Name | Kind | +| --- | --- | +| `dataset_path` | property | +| `max_host_memory_gb` | property | +| `num_threads` | property | + +### dataset_path + +```python +def dataset_path(self) +``` + +### max_host_memory_gb + +```python +def max_host_memory_gb(self) +``` + +### num_threads + +```python +def num_threads(self) +``` + ## build `@auto_sync_resources` @@ -385,6 +434,64 @@ version of cuVS is not guaranteed to work. ... "sqeuclidean") ``` +## materialize_to_hnswlib + +`@auto_sync_resources` + +```python +def materialize_to_hnswlib(MaterializeParams materialize_params, layered_artifact_path, output_path, dim, metric="sqeuclidean", resources=None) +``` + +Materialize a layered HNSW artifact into a standard hnswlib index file +on disk. + +Materializes a `GRAPH_ONLY` artifact (graph topology only, stored +in ACE order) plus a local dataset into a standard hnswlib index file, +without ever holding the full materialized index in host memory. The +resulting file is compatible with the original hnswlib library and can be +read back through `load()` with `hierarchy="cpu"`. The element data type +(float32, float16, uint8, int8) is inferred from the external dataset. +GRAPH_ONLY artifacts are currently produced through the C++ API. + +**Parameters** + +| Name | Type | Description | +| --- | --- | --- | +| `materialize_params` | `MaterializeParams` | Materialization parameters. `dataset_path` must point to the original-ID-ordered vectors used to build the artifact. | +| `layered_artifact_path` | `string` | Path to the layered HNSW artifact. | +| `output_path` | `string` | Path to the hnswlib index file to write. | +| `dim` | `int` | Dimensions of the training dataset. | +| `metric` | `string denoting the metric type, default="sqeuclidean"` | Valid values for metric: ["sqeuclidean", "inner_product"], where
- sqeuclidean is the euclidean distance without the square root operation, i.e.: distance(a,b) = \\sum_i (a_i - b_i)^2,
- inner_product distance is defined as distance(a, b) = \\sum_i a_i * b_i. | +| `resources` | `cuvs.common.Resources, optional` | | + +**Examples** + +```python +>>> import numpy as np +>>> from cuvs.neighbors import hnsw +>>> n_features = 50 +>>> # Assume a layered artifact was produced by an ACE GPU build and the +>>> # original-ID-ordered vectors are stored in "dataset.fbin". +>>> materialize_params = hnsw.MaterializeParams( +... dataset_path="dataset.fbin" +... ) +>>> hnsw.materialize_to_hnswlib( +... materialize_params, +... "layered_artifact.cuvs", +... "index.bin", +... n_features, +... metric="sqeuclidean", +... ) +>>> # The materialized index can be loaded as a standard hnswlib index. +>>> index = hnsw.load( +... hnsw.IndexParams(hierarchy="cpu"), +... "index.bin", +... n_features, +... np.float32, +... "sqeuclidean", +... ) +``` + ## save `@auto_sync_resources` diff --git a/java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java b/java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java index 50216729a8..4875ed55f5 100644 --- a/java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java +++ b/java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswIndex.java @@ -76,6 +76,45 @@ static HnswIndex build(CuVSResources resources, HnswIndexParams hnswParams, CuVS return CuVSProvider.provider().hnswIndexBuild(resources, hnswParams, dataset); } + /** + * Materializes a layered HNSW artifact into a standard hnswlib index file on + * disk. + * + * Materializes a {@code GRAPH_ONLY} artifact (graph topology only, + * stored in ACE order) plus a local dataset into a standard hnswlib index file, + * without ever holding the full materialized index in host memory. The + * resulting file is compatible with the original hnswlib library and can be read + * back with {@code hierarchy == CPU}. The element data type is inferred from the + * external dataset. GRAPH_ONLY artifacts are currently produced through the C++ + * API. + * + * @param resources The CuVS resources + * @param materializeParams Materialization parameters (dataset path, host-memory + * budget, threads) + * @param layeredArtifactPath Path to the layered HNSW artifact + * @param outputPath Path to the hnswlib index file to write + * @param dim The dimension of the vectors in the index + * @param metric The distance metric used to build the index + * @throws Throwable if an error occurs during materialization + */ + static void materializeToHnswlib( + CuVSResources resources, + HnswMaterializeParams materializeParams, + String layeredArtifactPath, + String outputPath, + int dim, + HnswIndexParams.CuvsDistanceType metric) + throws Throwable { + Objects.requireNonNull(resources); + Objects.requireNonNull(materializeParams); + Objects.requireNonNull(layeredArtifactPath); + Objects.requireNonNull(outputPath); + Objects.requireNonNull(metric); + CuVSProvider.provider() + .hnswMaterializeToHnswlib( + resources, materializeParams, layeredArtifactPath, outputPath, dim, metric); + } + /** * Builder helps configure and create an instance of {@link HnswIndex}. */ diff --git a/java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java b/java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java new file mode 100644 index 0000000000..3dd5fb3b33 --- /dev/null +++ b/java/cuvs-java/src/main/java/com/nvidia/cuvs/HnswMaterializeParams.java @@ -0,0 +1,134 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +package com.nvidia.cuvs; + +/** + * Parameters for materializing a layered HNSW artifact into a standard hnswlib + * index file on disk. + * + * @since 26.12 + */ +public class HnswMaterializeParams { + + private final String datasetPath; + private final double maxHostMemoryGb; + private final int numThreads; + + private HnswMaterializeParams(String datasetPath, double maxHostMemoryGb, int numThreads) { + this.datasetPath = datasetPath; + this.maxHostMemoryGb = maxHostMemoryGb; + this.numThreads = numThreads; + } + + /** + * Gets the local dataset path holding the original-ID-ordered vectors used to + * build the artifact. + * + * @return the dataset path, or null if not set + */ + public String getDatasetPath() { + return datasetPath; + } + + /** + * Gets the upper bound on host memory (in GiB) used for the base-topology + * reorder buffer. When {@code <= 0}, the whole base topology is reordered in a + * single in-memory pass. + * + * @return the max host memory in GiB (0 means a single in-memory pass) + */ + public double getMaxHostMemoryGb() { + return maxHostMemoryGb; + } + + /** + * Gets the number of host threads to use. When 0, the maximum number of + * threads is used. + * + * @return the number of threads + */ + public int getNumThreads() { + return numThreads; + } + + @Override + public String toString() { + return "HnswMaterializeParams [datasetPath=" + + datasetPath + + ", maxHostMemoryGb=" + + maxHostMemoryGb + + ", numThreads=" + + numThreads + + "]"; + } + + /** + * Builder configures and creates an instance of {@link HnswMaterializeParams}. + */ + public static class Builder { + + private String datasetPath; + private double maxHostMemoryGb = 0; + private int numThreads = 0; + + /** + * Constructs this Builder. + */ + public Builder() {} + + /** + * Sets the local dataset path holding the original-ID-ordered vectors used to + * build the artifact. Supported formats match layered deserialization: + * {@code .npy} and ANN benchmark {@code *.bin} files with a + * {@code [uint32 rows, uint32 cols]} header ({@code .fbin}, {@code .f16bin}, + * {@code .u8bin}, {@code .i8bin}). + * + * @param datasetPath the local dataset path + * @return an instance of Builder + */ + public Builder withDatasetPath(String datasetPath) { + this.datasetPath = datasetPath; + return this; + } + + /** + * Sets the upper bound on host memory (in GiB) used for the base-topology + * reorder buffer. + * + * When {@code <= 0} (default), the whole base topology is reordered in a single + * in-memory pass (no temporary files). When set, the base topology is reordered + * through bucketed temporary files so that peak host memory stays close to this + * budget. + * + * @param maxHostMemoryGb the max host memory in GiB + * @return an instance of Builder + */ + public Builder withMaxHostMemoryGb(double maxHostMemoryGb) { + this.maxHostMemoryGb = maxHostMemoryGb; + return this; + } + + /** + * Sets the number of host threads to use. When 0 (default), the maximum number + * of threads is used. + * + * @param numThreads the number of threads + * @return an instance of Builder + */ + public Builder withNumThreads(int numThreads) { + this.numThreads = numThreads; + return this; + } + + /** + * Builds an instance of {@link HnswMaterializeParams}. + * + * @return an instance of {@link HnswMaterializeParams} + */ + public HnswMaterializeParams build() { + return new HnswMaterializeParams(datasetPath, maxHostMemoryGb, numThreads); + } + } +} diff --git a/java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java b/java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java index fcd3481749..de6bcfbb17 100644 --- a/java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java +++ b/java/cuvs-java/src/main/java/com/nvidia/cuvs/spi/CuVSProvider.java @@ -166,6 +166,28 @@ HnswIndex.Builder newHnswIndexBuilder(CuVSResources cuVSResources) HnswIndex hnswIndexBuild(CuVSResources resources, HnswIndexParams hnswParams, CuVSMatrix dataset) throws Throwable; + /** + * Materializes a layered HNSW artifact into a standard hnswlib index file on disk. + * + * @param resources The CuVS resources + * @param materializeParams Materialization parameters (dataset path, host-memory budget, threads) + * @param layeredArtifactPath Path to the layered HNSW artifact + * @param outputPath Path to the hnswlib index file to write + * @param dim The dimension of the vectors in the index + * @param metric The distance metric used to build the index + * @throws Throwable if an error occurs during materialization + */ + default void hnswMaterializeToHnswlib( + CuVSResources resources, + HnswMaterializeParams materializeParams, + String layeredArtifactPath, + String outputPath, + int dim, + HnswIndexParams.CuvsDistanceType metric) + throws Throwable { + throw new UnsupportedOperationException(); + } + /** Creates a new TieredIndex Builder. */ TieredIndex.Builder newTieredIndexBuilder(CuVSResources cuVSResources) throws UnsupportedOperationException; diff --git a/java/cuvs-java/src/main/java22/com/nvidia/cuvs/internal/CuVSParamsHelper.java b/java/cuvs-java/src/main/java22/com/nvidia/cuvs/internal/CuVSParamsHelper.java index dd117515f3..ddda9a41dc 100644 --- a/java/cuvs-java/src/main/java22/com/nvidia/cuvs/internal/CuVSParamsHelper.java +++ b/java/cuvs-java/src/main/java22/com/nvidia/cuvs/internal/CuVSParamsHelper.java @@ -206,6 +206,27 @@ public void close() { } } + static CloseableHandle createHnswMaterializeParamsNative() { + try (var localArena = Arena.ofConfined()) { + var paramsPtrPtr = localArena.allocate(cuvsHnswMaterializeParams_t); + checkCuVSError( + cuvsHnswMaterializeParamsCreate(paramsPtrPtr), "cuvsHnswMaterializeParamsCreate"); + var paramsPtr = paramsPtrPtr.get(cuvsHnswMaterializeParams_t, 0L); + return new CloseableHandle() { + @Override + public MemorySegment handle() { + return paramsPtr; + } + + @Override + public void close() { + checkCuVSError( + cuvsHnswMaterializeParamsDestroy(paramsPtr), "cuvsHnswMaterializeParamsDestroy"); + } + }; + } + } + static CloseableHandle createTieredIndexParams() { try (var localArena = Arena.ofConfined()) { var paramsPtrPtr = localArena.allocate(cuvsTieredIndexParams_t); diff --git a/java/cuvs-java/src/main/java22/com/nvidia/cuvs/internal/HnswIndexImpl.java b/java/cuvs-java/src/main/java22/com/nvidia/cuvs/internal/HnswIndexImpl.java index 6f32e8deab..bda5a3611d 100644 --- a/java/cuvs-java/src/main/java22/com/nvidia/cuvs/internal/HnswIndexImpl.java +++ b/java/cuvs-java/src/main/java22/com/nvidia/cuvs/internal/HnswIndexImpl.java @@ -6,6 +6,7 @@ import static com.nvidia.cuvs.internal.CuVSParamsHelper.createHnswAceParamsNative; import static com.nvidia.cuvs.internal.CuVSParamsHelper.createHnswIndexParams; +import static com.nvidia.cuvs.internal.CuVSParamsHelper.createHnswMaterializeParamsNative; import static com.nvidia.cuvs.internal.common.LinkerHelper.C_FLOAT; import static com.nvidia.cuvs.internal.common.LinkerHelper.C_LONG; import static com.nvidia.cuvs.internal.common.Util.buildMemorySegment; @@ -19,6 +20,7 @@ import com.nvidia.cuvs.HnswAceParams; import com.nvidia.cuvs.HnswIndex; import com.nvidia.cuvs.HnswIndexParams; +import com.nvidia.cuvs.HnswMaterializeParams; import com.nvidia.cuvs.HnswQuery; import com.nvidia.cuvs.HnswSearchParams; import com.nvidia.cuvs.SearchResults; @@ -27,6 +29,7 @@ import com.nvidia.cuvs.internal.panama.cuvsHnswAceParams; import com.nvidia.cuvs.internal.panama.cuvsHnswIndex; import com.nvidia.cuvs.internal.panama.cuvsHnswIndexParams; +import com.nvidia.cuvs.internal.panama.cuvsHnswMaterializeParams; import com.nvidia.cuvs.internal.panama.cuvsHnswSearchParams; import java.io.InputStream; import java.lang.foreign.Arena; @@ -218,7 +221,7 @@ private IndexReference deserialize(InputStream inputStream) throws Throwable { } /** - * Allocates the configured search parameters in the MemorySegment. + * Allocates the configured index parameters in the MemorySegment. */ private CloseableHandle segmentFromIndexParams(HnswIndexParams params) { var hnswParams = createHnswIndexParams(); @@ -287,6 +290,56 @@ public static HnswIndex build( return new HnswIndexImpl(new IndexReference(hnswIndex), resources, hnswParams); } + /** + * Materializes a layered HNSW artifact into a standard hnswlib index file on disk. + * + * @param resources the CuVS resources + * @param materializeParams the materialization parameters + * @param layeredArtifactPath path to the layered HNSW artifact + * @param outputPath path to the hnswlib index file to write + * @param dim the dimension of the vectors in the index + * @param metric the distance metric used to build the index + * @throws Throwable if an error occurs during materialization + */ + public static void materializeToHnswlib( + CuVSResources resources, + HnswMaterializeParams materializeParams, + String layeredArtifactPath, + String outputPath, + int dim, + HnswIndexParams.CuvsDistanceType metric) + throws Throwable { + Objects.requireNonNull(resources); + Objects.requireNonNull(materializeParams); + Objects.requireNonNull(layeredArtifactPath); + Objects.requireNonNull(outputPath); + Objects.requireNonNull(metric); + + try (var localArena = Arena.ofConfined(); + var paramsHandle = createHnswMaterializeParamsNative()) { + MemorySegment paramsSeg = paramsHandle.handle(); + + String datasetPath = materializeParams.getDatasetPath(); + if (datasetPath != null) { + cuvsHnswMaterializeParams.dataset_path(paramsSeg, localArena.allocateFrom(datasetPath)); + } + cuvsHnswMaterializeParams.max_host_memory_gb( + paramsSeg, materializeParams.getMaxHostMemoryGb()); + cuvsHnswMaterializeParams.num_threads(paramsSeg, materializeParams.getNumThreads()); + + MemorySegment artifactSeg = buildMemorySegment(localArena, layeredArtifactPath); + MemorySegment outputSeg = buildMemorySegment(localArena, outputPath); + + try (var resourcesAccessor = resources.access()) { + var cuvsRes = resourcesAccessor.handle(); + int returnValue = + cuvsHnswMaterializeToHnswlib( + cuvsRes, paramsSeg, artifactSeg, outputSeg, dim, metric.value); + checkCuVSError(returnValue, "cuvsHnswMaterializeToHnswlib"); + } + } + } + private static CloseableHandle createHnswIndexParamsForBuild( Arena arena, HnswIndexParams params) { var hnswParams = createHnswIndexParams(); diff --git a/java/cuvs-java/src/main/java22/com/nvidia/cuvs/spi/JDKProvider.java b/java/cuvs-java/src/main/java22/com/nvidia/cuvs/spi/JDKProvider.java index bb27f36250..48f498d361 100644 --- a/java/cuvs-java/src/main/java22/com/nvidia/cuvs/spi/JDKProvider.java +++ b/java/cuvs-java/src/main/java22/com/nvidia/cuvs/spi/JDKProvider.java @@ -302,6 +302,19 @@ public HnswIndex hnswIndexBuild( return HnswIndexImpl.build(resources, hnswParams, dataset); } + @Override + public void hnswMaterializeToHnswlib( + CuVSResources resources, + HnswMaterializeParams materializeParams, + String layeredArtifactPath, + String outputPath, + int dim, + HnswIndexParams.CuvsDistanceType metric) + throws Throwable { + HnswIndexImpl.materializeToHnswlib( + resources, materializeParams, layeredArtifactPath, outputPath, dim, metric); + } + @Override public TieredIndex.Builder newTieredIndexBuilder(CuVSResources cuVSResources) { return TieredIndexImpl.newBuilder(Objects.requireNonNull(cuVSResources)); diff --git a/python/cuvs/cuvs/neighbors/hnsw/__init__.py b/python/cuvs/cuvs/neighbors/hnsw/__init__.py index f91835b7c5..cbac88163f 100644 --- a/python/cuvs/cuvs/neighbors/hnsw/__init__.py +++ b/python/cuvs/cuvs/neighbors/hnsw/__init__.py @@ -1,4 +1,4 @@ -# 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 @@ -7,11 +7,13 @@ ExtendParams, Index, IndexParams, + MaterializeParams, SearchParams, build, extend, from_cagra, load, + materialize_to_hnswlib, save, search, ) @@ -21,10 +23,12 @@ "IndexParams", "Index", "ExtendParams", + "MaterializeParams", "build", "extend", "SearchParams", "load", + "materialize_to_hnswlib", "save", "search", "from_cagra", diff --git a/python/cuvs/cuvs/neighbors/hnsw/hnsw.pxd b/python/cuvs/cuvs/neighbors/hnsw/hnsw.pxd index 9ffb295ad3..81abdf208a 100644 --- a/python/cuvs/cuvs/neighbors/hnsw/hnsw.pxd +++ b/python/cuvs/cuvs/neighbors/hnsw/hnsw.pxd @@ -1,5 +1,5 @@ # -# 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 # # cython: language_level=3 @@ -105,3 +105,24 @@ cdef extern from "cuvs/neighbors/hnsw.h" nogil: int32_t dim, cuvsDistanceType metric, cuvsHnswIndex_t index) except + + + ctypedef struct cuvsHnswMaterializeParams: + const char* dataset_path + double max_host_memory_gb + int num_threads + + ctypedef cuvsHnswMaterializeParams* cuvsHnswMaterializeParams_t + + cuvsError_t cuvsHnswMaterializeParamsCreate( + cuvsHnswMaterializeParams_t* params) + + cuvsError_t cuvsHnswMaterializeParamsDestroy( + cuvsHnswMaterializeParams_t params) + + cuvsError_t cuvsHnswMaterializeToHnswlib( + cuvsResources_t res, + cuvsHnswMaterializeParams_t params, + const char* layered_artifact_path, + const char* output_path, + int32_t dim, + cuvsDistanceType metric) except + diff --git a/python/cuvs/cuvs/neighbors/hnsw/hnsw.pyx b/python/cuvs/cuvs/neighbors/hnsw/hnsw.pyx index 91256a0e18..f7b0a8b807 100644 --- a/python/cuvs/cuvs/neighbors/hnsw/hnsw.pyx +++ b/python/cuvs/cuvs/neighbors/hnsw/hnsw.pyx @@ -213,7 +213,6 @@ cdef class IndexParams: def ace_params(self): return self._ace_params - cdef class Index: """ HNSW index object. This object stores the trained HNSW index state @@ -269,6 +268,68 @@ cdef class ExtendParams: return self.params.num_threads +cdef class MaterializeParams: + """ + Parameters for materializing a layered HNSW artifact into an hnswlib + index on disk. + + Parameters + ---------- + dataset_path : string, default = None (optional) + Local dataset path holding the original-ID-ordered vectors used to + build the artifact. Supported formats match layered deserialization: + `.npy` and ANN benchmark `*.bin` files with a + `[uint32 rows, uint32 cols]` header (`.fbin`, `.f16bin`, `.u8bin`, + `.i8bin`). + max_host_memory_gb : float, default = 0 (optional) + Upper bound on host memory (in GiB) used for the base-topology reorder + buffer. When <= 0, the whole base topology is reordered in a single + in-memory pass (no temporary files). When set, the base topology is + reordered through bucketed temporary files so that peak host memory + stays close to this budget. + num_threads : int, default = 0 (optional) + Number of host threads to use. When 0, the maximum number of threads + is used. + """ + + cdef cuvsHnswMaterializeParams* params + cdef object _dataset_path_bytes + + def __cinit__(self): + check_cuvs(cuvsHnswMaterializeParamsCreate(&self.params)) + self._dataset_path_bytes = None + + def __dealloc__(self): + if self.params is not NULL: + check_cuvs(cuvsHnswMaterializeParamsDestroy(self.params)) + + def __init__(self, *, + dataset_path=None, + max_host_memory_gb=0, + num_threads=0): + if dataset_path is not None: + self._dataset_path_bytes = dataset_path.encode('utf-8') + self.params.dataset_path = self._dataset_path_bytes + else: + self.params.dataset_path = NULL + self.params.max_host_memory_gb = max_host_memory_gb + self.params.num_threads = num_threads + + @property + def dataset_path(self): + if self.params.dataset_path is not NULL: + return self.params.dataset_path.decode('utf-8') + return None + + @property + def max_host_memory_gb(self): + return self.params.max_host_memory_gb + + @property + def num_threads(self): + return self.params.num_threads + + @auto_sync_resources def save(filename, Index index, resources=None): """ @@ -405,6 +466,85 @@ def load(IndexParams index_params, filename, dim, dtype, metric="sqeuclidean", return idx +@auto_sync_resources +def materialize_to_hnswlib(MaterializeParams materialize_params, + layered_artifact_path, + output_path, + dim, + metric="sqeuclidean", + resources=None): + """ + Materialize a layered HNSW artifact into a standard hnswlib index file + on disk. + + Materializes a `GRAPH_ONLY` artifact (graph topology only, stored + in ACE order) plus a local dataset into a standard hnswlib index file, + without ever holding the full materialized index in host memory. The + resulting file is compatible with the original hnswlib library and can be + read back through `load()` with `hierarchy="cpu"`. The element data type + (float32, float16, uint8, int8) is inferred from the external dataset. + GRAPH_ONLY artifacts are currently produced through the C++ API. + + Parameters + ---------- + materialize_params : MaterializeParams + Materialization parameters. `dataset_path` must point to the + original-ID-ordered vectors used to build the artifact. + layered_artifact_path : string + Path to the layered HNSW artifact. + output_path : string + Path to the hnswlib index file to write. + dim : int + Dimensions of the training dataset. + metric : string denoting the metric type, default="sqeuclidean" + Valid values for metric: ["sqeuclidean", "inner_product"], where + - sqeuclidean is the euclidean distance without the square root + operation, i.e.: distance(a,b) = \\sum_i (a_i - b_i)^2, + - inner_product distance is defined as + distance(a, b) = \\sum_i a_i * b_i. + {resources_docstring} + + Examples + -------- + >>> import numpy as np + >>> from cuvs.neighbors import hnsw + >>> n_features = 50 + >>> # Assume a layered artifact was produced by an ACE GPU build and the + >>> # original-ID-ordered vectors are stored in "dataset.fbin". + >>> materialize_params = hnsw.MaterializeParams( + ... dataset_path="dataset.fbin" + ... ) + >>> hnsw.materialize_to_hnswlib( + ... materialize_params, + ... "layered_artifact.cuvs", + ... "index.bin", + ... n_features, + ... metric="sqeuclidean", + ... ) + >>> # The materialized index can be loaded as a standard hnswlib index. + >>> index = hnsw.load( + ... hnsw.IndexParams(hierarchy="cpu"), + ... "index.bin", + ... n_features, + ... np.float32, + ... "sqeuclidean", + ... ) + """ + cdef string c_artifact = layered_artifact_path.encode('utf-8') + cdef string c_output = output_path.encode('utf-8') + cdef cuvsDistanceType distance_type = DISTANCE_TYPES[metric] + cdef cuvsResources_t res = resources.get_c_obj() + + check_cuvs(cuvsHnswMaterializeToHnswlib( + res, + materialize_params.params, + c_artifact.c_str(), + c_output.c_str(), + dim, + distance_type + )) + + @auto_sync_resources def from_cagra(IndexParams index_params, cagra.Index cagra_index, temporary_index_path=None, resources=None):