From 00ff5460a63db5445f53253804abd441ca1679eb Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Wed, 23 Sep 2026 11:02:39 +0200 Subject: [PATCH 01/20] Add general HNSW GRAPH_ONLY support --- cpp/CMakeLists.txt | 1 + cpp/include/cuvs/neighbors/hnsw.hpp | 115 +++++- cpp/src/neighbors/detail/hnsw.hpp | 384 +++++++++++++++--- .../neighbors/detail/hnsw_layered_gather.hpp | 21 + cpp/src/neighbors/hnsw_layered_gather.cu | 67 +++ .../neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu | 55 ++- cpp/src/util/serialize_validation.hpp | 8 + cpp/tests/CMakeLists.txt | 2 +- cpp/tests/neighbors/ann_hnsw_ace.cuh | 118 +++++- .../ann_hnsw_ace/test_float_uint32_t.cu | 64 +++ .../ann_hnsw_ace/test_half_uint32_t.cu | 48 +++ .../ann_hnsw_ace/test_int8_t_uint32_t.cu | 48 +++ .../ann_hnsw_ace/test_uint8_t_uint32_t.cu | 48 +++ .../ann_ivf_rabitq/test_float_int64_t.cu | 91 +++++ cpp/tests/util/serialize_validation_test.cpp | 44 ++ examples/cpp/src/hnsw_ace_layered_example.cu | 69 ++-- fern/pages/cpp_api/cpp-api-neighbors-hnsw.md | 60 ++- 17 files changed, 1094 insertions(+), 149 deletions(-) create mode 100644 cpp/src/neighbors/detail/hnsw_layered_gather.hpp create mode 100644 cpp/src/neighbors/hnsw_layered_gather.cu create mode 100644 cpp/tests/util/serialize_validation_test.cpp diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index ef9977f72c..4d7e11e876 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -1476,6 +1476,7 @@ if(NOT BUILD_CPU_ONLY) src/neighbors/composite/index.cu $<$:src/neighbors/cagra.cpp> $<$:src/neighbors/hnsw.cpp> + $<$:src/neighbors/hnsw_layered_gather.cu> src/neighbors/ivf_common.cu src/neighbors/ivf_flat_index.cpp ${ivf_flat_build_extend_inst_files} diff --git a/cpp/include/cuvs/neighbors/hnsw.hpp b/cpp/include/cuvs/neighbors/hnsw.hpp index fe412e8cc1..a50ba7ef48 100644 --- a/cpp/include/cuvs/neighbors/hnsw.hpp +++ b/cpp/include/cuvs/neighbors/hnsw.hpp @@ -45,9 +45,13 @@ enum class HnswHierarchy { NONE, CPU, GPU }; * @brief Output artifact format for an HNSW index * * `HNSWLIB` produces the standard hnswlib index format. `GRAPH_ONLY` stores the graph - * separately from vectors. The current implementation requires `HnswHierarchy::GPU` and - * disk-backed ACE. Load a graph-only artifact with the two-filename `deserialize` overload, which - * reads the vectors from a separate local dataset. + * separately from vectors and requires `HnswHierarchy::GPU`. Regular CAGRA, in-memory ACE, and + * disk-backed ACE all support this format. Load a graph-only artifact with the two-filename + * `deserialize` overload and a separate dataset in original row order before searching. + * + * In-memory builds return a file-backed handle that owns a temporary artifact. Call `serialize` + * before destroying that handle to preserve the artifact. Disk-backed ACE artifacts remain in + * the caller's build directory after the handle is destroyed. */ enum class HnswOutputFormat { HNSWLIB, @@ -64,8 +68,7 @@ enum class HnswOutputFormat { struct index_params : cuvs::neighbors::index_params { /** Hierarchy build type for HNSW index when converting from CAGRA index */ HnswHierarchy hierarchy = HnswHierarchy::GPU; - /** Output artifact format. Graph-only output currently requires a GPU hierarchy and disk-backed - * ACE. */ + /** Output artifact format. Graph-only output requires a GPU hierarchy. */ HnswOutputFormat output_format = HnswOutputFormat::HNSWLIB; /** Size of the candidate list during hierarchy construction when hierarchy is `CPU`*/ int ef_construction = 200; @@ -199,6 +202,10 @@ struct index : cuvs::neighbors::index { /** @brief Get file path for disk-backed index + + Graph-only artifacts from in-memory builds are temporary and remain valid while this handle is + alive. Use `serialize` to save a persistent copy. Disk-backed ACE artifacts remain in the + caller's build directory. */ virtual std::string file_path() const { return ""; } @@ -245,6 +252,11 @@ struct extend_params { * does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is * important to set `ace_params.build_dir` to a fast disk with sufficient storage size. * + * `GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. + * The returned handle references a graph artifact. Save it with `serialize`, then load it with + * the two-filename `deserialize` overload and an original-order dataset before searching. + * Temporary artifacts from in-memory builds are removed when the handle is destroyed. + * * NOTE: This function requires CUDA headers to be available at compile time. * * @param[in] res raft resources @@ -303,6 +315,11 @@ std::unique_ptr> build( * does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is * important to set `ace_params.build_dir` to a fast disk with sufficient storage size. * + * `GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. + * The returned handle references a graph artifact. Save it with `serialize`, then load it with + * the two-filename `deserialize` overload and an original-order dataset before searching. + * Temporary artifacts from in-memory builds are removed when the handle is destroyed. + * * NOTE: This function requires CUDA headers to be available at compile time. * * @param[in] res raft resources @@ -361,6 +378,11 @@ std::unique_ptr> build( * does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is * important to set `ace_params.build_dir` to a fast disk with sufficient storage size. * + * `GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. + * The returned handle references a graph artifact. Save it with `serialize`, then load it with + * the two-filename `deserialize` overload and an original-order dataset before searching. + * Temporary artifacts from in-memory builds are removed when the handle is destroyed. + * * NOTE: This function requires CUDA headers to be available at compile time. * * @param[in] res raft resources @@ -419,6 +441,11 @@ std::unique_ptr> build( * does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is * important to set `ace_params.build_dir` to a fast disk with sufficient storage size. * + * `GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. + * The returned handle references a graph artifact. Save it with `serialize`, then load it with + * the two-filename `deserialize` overload and an original-order dataset before searching. + * Temporary artifacts from in-memory builds are removed when the handle is destroyed. + * * NOTE: This function requires CUDA headers to be available at compile time. * * @param[in] res raft resources @@ -486,12 +513,18 @@ std::unique_ptr> build( * 3. `GPU`: The hierarchy is constructed on the GPU. * When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. * Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the - * local dataset. + * local dataset in original row order. In-memory conversion returns a handle owning a temporary + * artifact. Call `serialize` before destroying the handle to preserve it. + * + * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` + * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * * @param[in] res raft resources * @param[in] params hnsw index parameters * @param[in] cagra_index cagra index - * @param[in] dataset optional dataset to avoid extra memory copy when hierarchy is `CPU` + * @param[in] dataset optional original-order host vectors. For a GPU hierarchy, including + * `GRAPH_ONLY`, provide these if the CAGRA index has no attached device dataset. Rows and logical + * dimensions must match the CAGRA index. Attached padded device datasets are supported. * * Usage example: * @code{.cpp} @@ -526,12 +559,18 @@ std::unique_ptr> from_cagra( * 3. `GPU`: The hierarchy is constructed on the GPU. * When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. * Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the - * local dataset. + * local dataset in original row order. In-memory conversion returns a handle owning a temporary + * artifact. Call `serialize` before destroying the handle to preserve it. + * + * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` + * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * * @param[in] res raft resources * @param[in] params hnsw index parameters * @param[in] cagra_index cagra index - * @param[in] dataset optional dataset to avoid extra memory copy when hierarchy is `CPU` + * @param[in] dataset optional original-order host vectors. For a GPU hierarchy, including + * `GRAPH_ONLY`, provide these if the CAGRA index has no attached device dataset. Rows and logical + * dimensions must match the CAGRA index. Attached padded device datasets are supported. * * Usage example: * @code{.cpp} @@ -566,12 +605,18 @@ std::unique_ptr> from_cagra( * 3. `GPU`: The hierarchy is constructed on the GPU. * When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. * Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the - * local dataset. + * local dataset in original row order. In-memory conversion returns a handle owning a temporary + * artifact. Call `serialize` before destroying the handle to preserve it. + * + * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` + * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * * @param[in] res raft resources * @param[in] params hnsw index parameters * @param[in] cagra_index cagra index - * @param[in] dataset optional dataset to avoid extra memory copy when hierarchy is `CPU` + * @param[in] dataset optional original-order host vectors. For a GPU hierarchy, including + * `GRAPH_ONLY`, provide these if the CAGRA index has no attached device dataset. Rows and logical + * dimensions must match the CAGRA index. Attached padded device datasets are supported. * * Usage example: * @code{.cpp} @@ -606,12 +651,18 @@ std::unique_ptr> from_cagra( * 3. `GPU`: The hierarchy is constructed on the GPU. * When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. * Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the - * local dataset. + * local dataset in original row order. In-memory conversion returns a handle owning a temporary + * artifact. Call `serialize` before destroying the handle to preserve it. + * + * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` + * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * * @param[in] res raft resources * @param[in] params hnsw index parameters * @param[in] cagra_index cagra index - * @param[in] dataset optional dataset to avoid extra memory copy when hierarchy is `CPU` + * @param[in] dataset optional original-order host vectors. For a GPU hierarchy, including + * `GRAPH_ONLY`, provide these if the CAGRA index has no attached device dataset. Rows and logical + * dimensions must match the CAGRA index. Attached padded device datasets are supported. * * Usage example: * @code{.cpp} @@ -638,7 +689,11 @@ std::unique_ptr> from_cagra( * @brief Construct an hnswlib index from a device-standard CAGRA index. * * When the index has an attached device dataset view, `dataset` may be omitted. Otherwise pass a - * host matrix with the vectors (same contract as `device_padded_index`). + * host matrix with the vectors (same contract as `device_padded_index`). `GRAPH_ONLY` requires + * a GPU hierarchy and follows the same dataset and temporary-artifact contracts. + * + * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` + * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. */ std::unique_ptr> from_cagra( raft::resources const& res, @@ -671,6 +726,12 @@ std::unique_ptr> from_cagra( /** * @brief Construct an hnswlib index from a host-built CAGRA index. * Requires `dataset` for in-memory indices — host builds do not store vectors in the index. + * This also applies to `GRAPH_ONLY`, which requires a GPU hierarchy. Supply vectors in original + * row order with matching row count and logical dimensions. Serialize the returned temporary + * graph artifact before destroying the handle. + * + * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` + * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. */ std::unique_ptr> from_cagra( raft::resources const& res, @@ -703,6 +764,12 @@ std::unique_ptr> from_cagra( /** * @brief Construct an hnswlib index from a host-built CAGRA index (standard dataset layout). * Requires `dataset` for in-memory indices — host builds do not store vectors in the index. + * This also applies to `GRAPH_ONLY`, which requires a GPU hierarchy. Supply vectors in original + * row order with matching row count and logical dimensions. Serialize the returned temporary + * graph artifact before destroying the handle. + * + * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` + * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. */ std::unique_ptr> from_cagra( raft::resources const& res, @@ -1103,7 +1170,10 @@ void search(raft::resources const& res, * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. - * Load it with the two-filename `deserialize` overload and a local dataset. + * Load it with the two-filename `deserialize` overload and a local dataset in original row order. + * Serialize to a filename outside the handle's temporary directory to preserve the artifact + * beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer + * ownership. * * @param[in] res raft resources * @param[in] filename path to the serialized HNSW output @@ -1134,7 +1204,10 @@ void serialize(raft::resources const& res, const std::string& filename, const in * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. - * Load it with the two-filename `deserialize` overload and a local dataset. + * Load it with the two-filename `deserialize` overload and a local dataset in original row order. + * Serialize to a filename outside the handle's temporary directory to preserve the artifact + * beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer + * ownership. * * @param[in] res raft resources * @param[in] filename path to the serialized HNSW output @@ -1165,7 +1238,10 @@ void serialize(raft::resources const& res, const std::string& filename, const in * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. - * Load it with the two-filename `deserialize` overload and a local dataset. + * Load it with the two-filename `deserialize` overload and a local dataset in original row order. + * Serialize to a filename outside the handle's temporary directory to preserve the artifact + * beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer + * ownership. * * @param[in] res raft resources * @param[in] filename path to the serialized HNSW output @@ -1196,7 +1272,10 @@ void serialize(raft::resources const& res, const std::string& filename, const in * However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib * library. * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. - * Load it with the two-filename `deserialize` overload and a local dataset. + * Load it with the two-filename `deserialize` overload and a local dataset in original row order. + * Serialize to a filename outside the handle's temporary directory to preserve the artifact + * beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer + * ownership. * * @param[in] res raft resources * @param[in] filename path to the serialized HNSW output diff --git a/cpp/src/neighbors/detail/hnsw.hpp b/cpp/src/neighbors/detail/hnsw.hpp index 19c0afd820..a77b5c6362 100644 --- a/cpp/src/neighbors/detail/hnsw.hpp +++ b/cpp/src/neighbors/detail/hnsw.hpp @@ -8,6 +8,7 @@ #include "../../core/nvtx.hpp" #include "../../core/omp_wrapper.hpp" #include "hnsw_layered_format.hpp" +#include "hnsw_layered_gather.hpp" #include #include @@ -141,6 +142,35 @@ class exclusive_hnsw_output_file { std::unique_ptr stream_; }; +// Own only directories allocated here; user-supplied ACE directories remain persistent. +class owned_hnsw_directory { + public: + owned_hnsw_directory() + { + auto pattern = (std::filesystem::temp_directory_path() / "cuvs-hnsw-XXXXXX").string(); + RAFT_EXPECTS(::mkdtemp(pattern.data()) != nullptr, + "Cannot create HNSW artifact directory: %s", + std::strerror(errno)); + try { + path_ = pattern; + } catch (...) { + (void)::rmdir(pattern.c_str()); + throw; + } + } + owned_hnsw_directory(const owned_hnsw_directory&) = delete; + owned_hnsw_directory& operator=(const owned_hnsw_directory&) = delete; + ~owned_hnsw_directory() + { + std::error_code error; + std::filesystem::remove_all(path_, error); + } + const std::filesystem::path& path() const { return path_; } + + private: + std::filesystem::path path_; +}; + template inline constexpr bool is_cagra_hnsw_export_index_v = std::is_same_v> || @@ -266,6 +296,10 @@ struct index_impl : index { appr_alg_ = std::move(index); } + void set_construction_dtype(layered_hnsw_dtype dtype) { construction_dtype_ = dtype; } + + auto construction_dtype() const -> layered_hnsw_dtype { return construction_dtype_; } + /** @brief Get space */ @@ -279,6 +313,11 @@ struct index_impl : index { */ void set_file_descriptor(cuvs::util::file_descriptor&& fd) { hnsw_fd_.emplace(std::move(fd)); } + void set_artifact_directory(std::unique_ptr directory) + { + artifact_directory_ = std::move(directory); + } + /** @brief Get file descriptor */ @@ -336,7 +375,10 @@ struct index_impl : index { private: mutable std::unique_ptr::type>> appr_alg_; std::unique_ptr::type>> space_; + // Declared before the descriptor so the file is closed before directory cleanup. + std::unique_ptr artifact_directory_; std::optional hnsw_fd_; + layered_hnsw_dtype construction_dtype_ = layered_hnsw_dtype::unknown; }; template @@ -1073,11 +1115,83 @@ void write_layered_base_links_from_disk(const CagraIndexT& index_, } template -auto serialize_to_layered_hnswlib_from_disk( +void write_layered_base_links_from_inmem(raft::resources const& res, + const CagraIndexT& index_, + cuvs::util::file_descriptor& output, + size_t nodes_offset, + size_t links_offset, + size_t link_row_bytes, + size_t maxM0) +{ + using IdxT = typename CagraIndexT::index_type; + auto graph = index_.graph(); + const auto rows = static_cast(index_.size()); + const auto degree = static_cast(graph.extent(1)); + RAFT_EXPECTS(static_cast(graph.extent(0)) == rows && graph.data_handle() != nullptr, + "Missing or mismatched in-memory CAGRA graph"); + RAFT_EXPECTS(degree > 0 && degree <= maxM0, "Invalid base graph degree"); + const auto batch_rows = + std::min(rows, + std::max( + 1, (64 * 1024 * 1024) / (degree * sizeof(IdxT) + sizeof(IdxT) + link_row_bytes))); + auto neighbors = raft::make_host_matrix(batch_rows, degree); + std::vector nodes(batch_rows); + std::vector links(batch_rows * link_row_bytes); + auto* graph_ptr = graph.data_handle(); + cudaPointerAttributes attributes; + RAFT_CUDA_TRY(cudaPointerGetAttributes(&attributes, graph_ptr)); + const bool host_accessible = + attributes.type == cudaMemoryTypeUnregistered || attributes.hostPointer != nullptr; + if (attributes.hostPointer != nullptr) { + graph_ptr = static_cast(attributes.hostPointer); + } + // Host-accessible graphs may have been produced asynchronously on the resource stream. + raft::resource::sync_stream(res); + for (size_t start = 0; start < rows; start += batch_rows) { + const auto count = std::min(batch_rows, rows - start); + if (host_accessible) { + for (size_t row = 0; row < count; ++row) { + std::copy_n(graph_ptr + (start + row) * graph.stride(0), + degree, + neighbors.data_handle() + row * degree); + } + } else { + raft::copy_matrix(neighbors.data_handle(), + degree, + graph_ptr + start * graph.stride(0), + graph.stride(0), + degree, + count, + raft::resource::get_cuda_stream(res)); + raft::resource::sync_stream(res); + } + std::fill(links.begin(), links.end(), 0); + for (size_t row = 0; row < count; ++row) { + nodes[row] = static_cast(start + row); + auto* destination = links.data() + row * link_row_bytes; + const auto length = static_cast(degree); + std::memcpy(destination, &length, sizeof(length)); + for (size_t col = 0; col < degree; ++col) { + RAFT_EXPECTS(neighbors(row, col) < rows, "Invalid neighbor id in CAGRA graph"); + } + std::memcpy(destination + sizeof(length), + neighbors.data_handle() + row * degree, + degree * sizeof(IdxT)); + } + cuvs::util::write_large_file( + output, nodes.data(), count * sizeof(IdxT), nodes_offset + start * sizeof(IdxT)); + cuvs::util::write_large_file( + output, links.data(), count * link_row_bytes, links_offset + start * link_row_bytes); + } +} + +template +auto serialize_to_layered_hnswlib( raft::resources const& res, const cuvs::neighbors::hnsw::index_params& params, const CagraIndexT& index_, - raft::host_matrix_view dataset) -> std::string + std::optional> dataset, + const std::filesystem::path& artifact_file) -> std::string { using IdxT = typename CagraIndexT::index_type; static_assert(is_cagra_hnsw_export_index_v, @@ -1085,30 +1199,39 @@ auto serialize_to_layered_hnswlib_from_disk( raft::common::nvtx::range fun_scope("hnsw::serialize_layered"); const auto total_start_time = std::chrono::steady_clock::now(); - RAFT_EXPECTS(index_.graph_fd().has_value() && index_.mapping_fd().has_value(), - "Layered HNSW serialization requires a disk-backed ACE graph and mapping."); - RAFT_EXPECTS(static_cast(dataset.extent(0)) == static_cast(index_.size()), - "Dataset rows (%zu) must match index size (%zu)", - static_cast(dataset.extent(0)), - static_cast(index_.size())); - RAFT_EXPECTS(static_cast(dataset.extent(1)) == static_cast(index_.dim()), - "Dataset cols (%zu) must match index dimensions (%zu)", - static_cast(dataset.extent(1)), - static_cast(index_.dim())); - - const auto graph_path = index_.graph_fd()->get_path(); - RAFT_EXPECTS(!graph_path.empty(), "Unable to get path from graph file descriptor"); - const auto ace_dir = std::filesystem::path(graph_path).parent_path(); - const auto artifact_file = ace_dir / "hnsw_index.cuvs"; - std::filesystem::create_directories(ace_dir); + const bool disk_graph = index_.graph_fd().has_value(); + const T* source_dataset = nullptr; + size_t source_stride = 0; + bool device_dataset = false; + size_t dataset_rows = 0, dataset_dim = 0; + if (dataset.has_value()) { + source_dataset = dataset->data_handle(); + source_stride = dataset->stride(0); + dataset_rows = dataset->extent(0); + dataset_dim = dataset->extent(1); + } else if constexpr (is_device_cagra_hnsw_export_index_v) { + RAFT_EXPECTS(!disk_graph, "Layered HNSW requires the original-order host dataset for disk ACE"); + const auto view = index_.dataset(); + source_dataset = view.view().data_handle(); + source_stride = view.stride(); + dataset_rows = view.n_rows(); + dataset_dim = view.dim(); + device_dataset = true; + } + RAFT_EXPECTS(source_dataset != nullptr, "Layered HNSW requires a dataset"); + RAFT_EXPECTS(dataset_rows == static_cast(index_.size()), + "Dataset rows must match index size"); + RAFT_EXPECTS(index_.dim() == 0 || dataset_dim == static_cast(index_.dim()), + "Dataset columns must match index dimensions"); + RAFT_EXPECTS(dataset_rows > 0 && dataset_dim > 0 && index_.graph_degree() > 0, + "Layered HNSW requires nonempty vectors and graph"); auto n_rows = static_cast(index_.size()); - auto dim = static_cast(index_.dim()); + auto dim = dataset_dim; auto graph_degree_int = static_cast(index_.graph_degree()); - auto hnsw_index = - std::make_unique>(index_.dim(), index_.metric(), params.hierarchy); - auto appr_algo = std::make_unique::type>>( + auto hnsw_index = std::make_unique>(dim, index_.metric(), params.hierarchy); + auto appr_algo = std::make_unique::type>>( hnsw_index->get_space(), 1, (graph_degree_int + 1) / 2, params.ef_construction); RAFT_LOG_INFO("Layered HNSW artifact: generating hierarchy levels"); @@ -1199,12 +1322,22 @@ auto serialize_to_layered_hnswlib_from_disk( RAFT_LOG_INFO("Layered HNSW artifact: writing hnswlib-ready base topology section"); const auto layer0_start_time = std::chrono::steady_clock::now(); - write_layered_base_links_from_disk(index_, - artifact_fd, - base_nodes_offset, - base_links_offset, - metadata.base_link_row_bytes, - metadata.maxM0); + if (disk_graph) { + write_layered_base_links_from_disk(index_, + artifact_fd, + base_nodes_offset, + base_links_offset, + metadata.base_link_row_bytes, + metadata.maxM0); + } else { + write_layered_base_links_from_inmem(res, + index_, + artifact_fd, + base_nodes_offset, + base_links_offset, + metadata.base_link_row_bytes, + metadata.maxM0); + } const auto layer0_elapsed_ms = elapsed_ms_since(layer0_start_time); static_cast(layer0_elapsed_ms); RAFT_LOG_INFO( @@ -1219,12 +1352,21 @@ auto serialize_to_layered_hnswlib_from_disk( const auto gather_start_time = std::chrono::steady_clock::now(); auto host_query_set = raft::make_host_matrix(static_cast(hierarchy.promoted_count()), dim); + if (device_dataset) { + gather_layered_hnsw_vectors(res, + source_dataset, + source_stride * sizeof(T), + dim * sizeof(T), + hierarchy.row_ids_by_level.data_handle() + hierarchy.hist[0], + hierarchy.promoted_count(), + host_query_set.data_handle()); + } else { + raft::resource::sync_stream(res); #pragma omp parallel for - for (int64_t i = 0; i < static_cast(n_rows); i++) { - if (hierarchy.levels(i) > 0) { - const auto query_row = hierarchy.position_by_row_id(i) - hierarchy.hist[0]; - auto* dst = host_query_set.data_handle() + query_row * dim; - std::copy(&dataset(i, 0), &dataset(i, 0) + dim, dst); + for (int64_t i = 0; i < static_cast(hierarchy.promoted_count()); ++i) { + const auto row = hierarchy.row_ids_by_level(hierarchy.hist[0] + i); + std::copy_n( + source_dataset + row * source_stride, dim, host_query_set.data_handle() + i * dim); } } const auto gather_elapsed_ms = elapsed_ms_since(gather_start_time); @@ -2232,6 +2374,32 @@ std::unique_ptr> from_cagra( } } + if (params.output_format == HnswOutputFormat::GRAPH_ONLY) { + RAFT_EXPECTS(!cagra_index.source_indices().has_value(), + "Layered HNSW does not support composite CAGRA source index mappings"); + auto hnsw_index = + std::make_unique>(dataset.has_value() ? dataset->extent(1) : cagra_index.dim(), + cagra_index.metric(), + params.hierarchy, + params.output_format); + std::unique_ptr directory; + std::filesystem::path artifact_path; + if (cagra_index.graph_fd().has_value()) { + RAFT_EXPECTS(cagra_index.graph_fd()->is_valid(), "Invalid disk graph descriptor"); + const auto graph_path = cagra_index.graph_fd()->get_path(); + RAFT_EXPECTS(!graph_path.empty(), "Unable to get graph file path"); + RAFT_EXPECTS(dataset.has_value(), "Layered HNSW requires the original-order host dataset"); + artifact_path = std::filesystem::path(graph_path).parent_path() / "hnsw_index.cuvs"; + } else { + directory = std::make_unique(); + artifact_path = directory->path() / "hnsw_index.cuvs"; + } + serialize_to_layered_hnswlib(res, params, cagra_index, dataset, artifact_path); + hnsw_index->set_file_descriptor(cuvs::util::file_descriptor(artifact_path.string(), O_RDONLY)); + hnsw_index->set_artifact_directory(std::move(directory)); + return hnsw_index; + } + // special treatment for index on disk if (cagra_index.dataset_fd().has_value() && cagra_index.graph_fd().has_value()) { // Get directory from graph file descriptor @@ -2247,18 +2415,6 @@ std::unique_ptr> from_cagra( std::filesystem::exists(index_directory) && std::filesystem::is_directory(index_directory), "Directory '%s' does not exist", index_directory.c_str()); - if (params.output_format == HnswOutputFormat::GRAPH_ONLY) { - RAFT_EXPECTS(dataset.has_value(), - "Layered HNSW serialization requires the original-order dataset."); - auto artifact_path = - serialize_to_layered_hnswlib_from_disk(res, params, cagra_index, dataset.value()); - - auto hnsw_index = std::make_unique>( - cagra_index.dim(), cagra_index.metric(), params.hierarchy, params.output_format); - hnsw_index->set_file_descriptor(cuvs::util::file_descriptor(artifact_path, O_RDONLY)); - return hnsw_index; - } - std::string index_filename = (std::filesystem::path(index_directory) / "hnsw_index.bin").string(); exclusive_hnsw_output_file output(index_filename); @@ -2278,14 +2434,8 @@ std::unique_ptr> from_cagra( return hnsw_index; } - RAFT_EXPECTS(params.output_format == HnswOutputFormat::HNSWLIB, - "GRAPH_ONLY requires disk-backed ACE build artifacts"); - - // In-memory CAGRA index: the resulting HNSW index might still not fit in host memory. - // Estimate its host footprint and, if it does not fit, spill it to disk via - // serialize_to_hnswlib_from_inmem instead of constructing it in RAM (NONE/GPU only; - // the CPU hierarchy is not supported by the batched serializer, and GRAPH_ONLY is only - // produced from disk-backed ACE artifacts handled above). + // In-memory HNSW may still exceed host memory; spill the vector-bearing format + // through the existing batched serializer when necessary (NONE/GPU only). if (params.hierarchy == HnswHierarchy::NONE || params.hierarchy == HnswHierarchy::GPU) { int64_t n_rows = dataset.has_value() ? dataset->extent(0) : cagra_index.size(); int64_t dim = dataset.has_value() ? dataset->extent(1) : cagra_index.dim(); @@ -2478,11 +2628,126 @@ void search(raft::resources const& res, } } +// Reconstruct the graph-only artifact from a loaded index without retaining its source file. +template +void serialize_loaded_layered_hnsw(const index_impl& idx, const std::string& filename) +{ + using algorithm_type = hnswlib::HierarchicalNSW::type>; + auto* graph = static_cast(idx.get_index()); + RAFT_EXPECTS(graph != nullptr, "Layered HNSW index has no loaded graph"); + + layered_hnsw_file_metadata metadata; + metadata.n_rows = graph->cur_element_count; + metadata.dim = idx.dim(); + metadata.M = graph->M_; + metadata.maxM = graph->maxM_; + metadata.maxM0 = graph->maxM0_; + metadata.ef_construction = graph->ef_construction_; + metadata.mult = graph->mult_; + metadata.maxlevel = graph->maxlevel_; + metadata.enterpoint_node = graph->enterpoint_node_; + metadata.levels_bytes = metadata.n_rows; + metadata.base_nodes_bytes = metadata.n_rows * sizeof(uint32_t); + metadata.base_link_row_bytes = graph->size_links_level0_; + metadata.base_links_bytes = metadata.n_rows * metadata.base_link_row_bytes; + metadata.upper_link_row_bytes = graph->size_links_per_element_; + + std::vector levels(metadata.n_rows); + for (size_t row = 0; row < metadata.n_rows; ++row) { + levels[row] = static_cast(graph->element_levels_[row]); + metadata.base_degree = std::max( + metadata.base_degree, static_cast(graph->getListCount(graph->get_linklist0(row)))); + } + std::vector descriptors; + for (int level = 1; level <= metadata.maxlevel; ++level) { + const auto count = static_cast(std::count_if( + levels.begin(), levels.end(), [level](uint8_t value) { return value >= level; })); + const auto degree = hnsw_upper_layer_degree(count, metadata.M); + metadata.layers.push_back({static_cast(level), + count, + degree, + metadata.upper_nodes_count, + metadata.upper_nodes_count}); + descriptors.push_back({static_cast(level), + count, + degree, + metadata.upper_nodes_count, + metadata.upper_nodes_count}); + metadata.upper_nodes_count += count; + } + metadata.upper_nodes_bytes = metadata.upper_nodes_count * sizeof(uint32_t); + metadata.upper_links_bytes = metadata.upper_nodes_count * metadata.upper_link_row_bytes; + auto header = make_layered_hnsw_header(metadata, idx.metric()); + header.construction_dtype = static_cast(idx.construction_dtype()); + const size_t descriptors_bytes = descriptors.size() * sizeof(layered_hnsw_layer_descriptor); + const size_t levels_offset = align_up(sizeof(header) + descriptors_bytes, layered_hnsw_alignment); + const size_t base_nodes_offset = levels_offset + metadata.levels_bytes; + const size_t base_links_offset = base_nodes_offset + metadata.base_nodes_bytes; + const size_t upper_nodes_offset = base_links_offset + metadata.base_links_bytes; + const size_t upper_links_offset = upper_nodes_offset + metadata.upper_nodes_bytes; + + const std::filesystem::path path(filename); + if (!path.parent_path().empty()) { std::filesystem::create_directories(path.parent_path()); } + exclusive_hnsw_temp_file output(path); + cuvs::util::file_descriptor fd(output.temporary_path(), O_RDWR); + cuvs::util::write_large_file(fd, &header, sizeof(header), 0); + if (descriptors_bytes > 0) { + cuvs::util::write_large_file(fd, descriptors.data(), descriptors_bytes, sizeof(header)); + } + cuvs::util::write_large_file(fd, levels.data(), levels.size(), levels_offset); + + // Bound scratch memory and batch writes instead of issuing one write per adjacency row. + const size_t batch_size = std::max(1, (64 * 1024 * 1024) / metadata.base_link_row_bytes); + std::vector nodes; + std::vector links; + nodes.reserve(std::min(batch_size, metadata.n_rows)); + links.reserve(std::min(batch_size, metadata.n_rows) * metadata.base_link_row_bytes); + auto write_layer = [&](int level, size_t node_offset, size_t link_offset, size_t row_bytes) { + size_t written = 0; + auto flush = [&] { + if (nodes.empty()) { return; } + cuvs::util::write_large_file(fd, + nodes.data(), + nodes.size() * sizeof(uint32_t), + node_offset + written * sizeof(uint32_t)); + cuvs::util::write_large_file( + fd, links.data(), links.size(), link_offset + written * row_bytes); + written += nodes.size(); + nodes.clear(); + links.clear(); + }; + for (size_t row = 0; row < metadata.n_rows; ++row) { + if (levels[row] < level) { continue; } + nodes.push_back(static_cast(row)); + const auto* link = reinterpret_cast( + level == 0 ? graph->get_linklist0(row) : graph->get_linklist(row, level)); + links.insert(links.end(), link, link + row_bytes); + if (nodes.size() == batch_size) { flush(); } + } + flush(); + }; + write_layer(0, base_nodes_offset, base_links_offset, metadata.base_link_row_bytes); + for (const auto& layer : metadata.layers) { + write_layer(static_cast(layer.level), + upper_nodes_offset + layer.node_offset * sizeof(uint32_t), + upper_links_offset + layer.link_offset * metadata.upper_link_row_bytes, + metadata.upper_link_row_bytes); + } + // Match serialize's overwrite semantics while publishing only complete artifacts. + std::filesystem::rename(output.temporary_path(), path); +} + template void serialize(raft::resources const& res, const std::string& filename, const index& idx) { auto* idx_impl = dynamic_cast*>(&idx); + if (idx_impl && idx_impl->output_format() == HnswOutputFormat::GRAPH_ONLY && + idx_impl->get_index() != nullptr) { + serialize_loaded_layered_hnsw(*idx_impl, filename); + return; + } + // Check if this is a disk-based index (created from disk-backed CAGRA) if (idx_impl && idx_impl->file_descriptor().has_value()) { // For disk-based indexes, copy the existing file to the new location @@ -2878,6 +3143,7 @@ auto deserialize_layered_hnswlib(raft::resources const& res, throughput_gib_per_s(upper_bytes_read, upper_read_time), elapsed_ms(upper_copy_time)); + hnsw_index->set_construction_dtype(construction_dtype); hnsw_index->set_index(std::move(appr_algo)); RAFT_LOG_INFO("Layered HNSW load: total deserialize completed in %ld ms", elapsed_ms_since(total_start_time)); @@ -2940,18 +3206,6 @@ std::unique_ptr> build(raft::resources const& res, common::nvtx::range fun_scope("hnsw::build"); validate_output_format(params); - // GRAPH_ONLY materializes the layered index from disk-backed ACE artifacts, so it - // requires a GPU hierarchy and ACE disk mode with a build directory. Validate up front; this - // also forces the CAGRA build below onto the ACE disk path (use_ace becomes true because ACE - // params are set). - if (params.output_format == HnswOutputFormat::GRAPH_ONLY) { - RAFT_EXPECTS(std::holds_alternative(params.graph_build_params), - "GRAPH_ONLY requires ACE parameters to be configured"); - const auto& ace = std::get(params.graph_build_params); - RAFT_EXPECTS(ace.use_disk, "GRAPH_ONLY requires ACE disk mode (ace_params.use_disk = true)"); - RAFT_EXPECTS(!ace.build_dir.empty(), "GRAPH_ONLY requires ace_params.build_dir to be set"); - } - cuvs::neighbors::cagra::index_params cagra_params = cagra::index_params::from_hnsw_params(dataset.extents(), params.M, diff --git a/cpp/src/neighbors/detail/hnsw_layered_gather.hpp b/cpp/src/neighbors/detail/hnsw_layered_gather.hpp new file mode 100644 index 0000000000..b3e491f1aa --- /dev/null +++ b/cpp/src/neighbors/detail/hnsw_layered_gather.hpp @@ -0,0 +1,21 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include +#include +#include + +namespace cuvs::neighbors::hnsw::detail { +// Gather only selected device rows to a compact host matrix using bounded device staging. +// Row IDs must be valid source rows; host_output must hold rows * row_bytes bytes. +void gather_layered_hnsw_vectors(raft::resources const& res, + const void* source, + size_t source_stride_bytes, + size_t row_bytes, + const size_t* host_row_ids, + size_t rows, + void* host_output); +} // namespace cuvs::neighbors::hnsw::detail diff --git a/cpp/src/neighbors/hnsw_layered_gather.cu b/cpp/src/neighbors/hnsw_layered_gather.cu new file mode 100644 index 0000000000..073da3304a --- /dev/null +++ b/cpp/src/neighbors/hnsw_layered_gather.cu @@ -0,0 +1,67 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +#include "detail/hnsw_layered_gather.hpp" + +#include +#include +#include +#include +#include + +namespace cuvs::neighbors::hnsw::detail { +namespace { +__global__ void gather_rows(const uint8_t* source, + size_t stride, + size_t row_bytes, + const size_t* row_ids, + size_t bytes, + uint8_t* output) +{ + for (size_t i = blockIdx.x * size_t(blockDim.x) + threadIdx.x; i < bytes; + i += size_t(blockDim.x) * gridDim.x) { + output[i] = source[size_t(row_ids[i / row_bytes]) * stride + i % row_bytes]; + } +} +} // namespace + +void gather_layered_hnsw_vectors(raft::resources const& res, + const void* source, + size_t source_stride_bytes, + size_t row_bytes, + const size_t* host_row_ids, + size_t rows, + void* host_output) +{ + if (rows == 0) { return; } + RAFT_EXPECTS(row_bytes > 0 && source_stride_bytes >= row_bytes, "Invalid vector row stride"); + const auto stream = raft::resource::get_cuda_stream(res); + const auto batch_rows = + std::min(rows, std::max(1, (64 * 1024 * 1024) / (row_bytes + sizeof(size_t)))); + rmm::device_uvector row_ids(batch_rows, stream); + rmm::device_uvector values(batch_rows * row_bytes, stream); + for (size_t start = 0; start < rows; start += batch_rows) { + const auto count = std::min(batch_rows, rows - start); + RAFT_CUDA_TRY(cudaMemcpyAsync(row_ids.data(), + host_row_ids + start, + count * sizeof(size_t), + cudaMemcpyHostToDevice, + stream)); + gather_rows<<(65535, (count * row_bytes + 255) / 256), 256, 0, stream>>>( + static_cast(source), + source_stride_bytes, + row_bytes, + row_ids.data(), + count * row_bytes, + values.data()); + RAFT_CUDA_TRY(cudaGetLastError()); + RAFT_CUDA_TRY(cudaMemcpyAsync(static_cast(host_output) + start * row_bytes, + values.data(), + count * row_bytes, + cudaMemcpyDeviceToHost, + stream)); + raft::resource::sync_stream(res); + } +} +} // namespace cuvs::neighbors::hnsw::detail diff --git a/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu b/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu index e27d34557a..76d1de87d8 100644 --- a/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu +++ b/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu @@ -191,8 +191,27 @@ void IVFGPU::load(const char* filename) // New change: host copy of ivf. AllocateHostMemory(); - read_into_device_host( - short_data_.data_handle(), short_data_host_.data_handle(), GetShortDataBytesSimple()); + // The persisted format is vector-major within each cluster. Search and quantization + // use word-major codes, so convert at the serialization boundary. + const size_t words_per_vector = DQ->short_code_length(); + std::vector disk_codes(max_cluster_length * words_per_vector); + size_t code_offset = 0; + for (auto cluster_size : cluster_sizes) { + const size_t cluster_words = cluster_size * words_per_vector; + read_exact(disk_codes.data(), cluster_words * sizeof(uint32_t)); + for (size_t word = 0; word < words_per_vector; ++word) { + for (size_t row = 0; row < cluster_size; ++row) { + short_data_host_(code_offset + word * cluster_size + row) = + disk_codes[row * words_per_vector + word]; + } + } + code_offset += cluster_words; + } + raft::copy(short_data_.data_handle(), + short_data_host_.data_handle(), + GetShortDataBytesSimple() / sizeof(uint32_t), + stream_); + raft::resource::sync_stream(handle_); read_into_device(short_factors_batch_.data_handle(), GetShortDataFactorBytesBatch()); // SoA layout - no copy on CPU read_into_device_host( @@ -322,17 +341,37 @@ void IVFGPU::save(const char* filename) const this->initializer->SaveCentroids(output, filename); // Compute sizes for device arrays. - size_t short_data_size = GetShortDataBytesSimple(); - size_t long_code_size = GetLongCodeBytes(); - size_t ex_factor_size = GetExFactorBytes(); - size_t ids_size = GetPIDsBytes(); + size_t long_code_size = GetLongCodeBytes(); + size_t ex_factor_size = GetExFactorBytes(); + size_t ids_size = GetPIDsBytes(); // for batch data size_t short_factors_size = GetShortDataFactorBytesBatch(); raft::resource::sync_stream(handle_); - // Write raw arrays to file. - output.write_device(short_data_.data_handle(), short_data_size); + // Preserve the legacy vector-major disk layout while keeping word-major codes + // in memory. Bound staging memory to the largest cluster. + const size_t words_per_vector = DQ->short_code_length(); + const size_t max_cluster_size = *std::max_element(cluster_sizes.begin(), cluster_sizes.end()); + std::vector memory_codes(max_cluster_size * words_per_vector); + std::vector disk_codes(max_cluster_size * words_per_vector); + for (auto const& cluster : h_cluster_meta) { + const size_t cluster_words = cluster.num * words_per_vector; + if (cluster_words == 0) { continue; } + raft::copy(memory_codes.data(), + short_data_.data_handle() + cluster.start_index * words_per_vector, + cluster_words, + stream_); + raft::resource::sync_stream(handle_); + for (size_t row = 0; row < cluster.num; ++row) { + for (size_t word = 0; word < words_per_vector; ++word) { + disk_codes[row * words_per_vector + word] = memory_codes[word * cluster.num + row]; + } + } + write_exact(disk_codes.data(), cluster_words * sizeof(uint32_t)); + } + + // The remaining arrays have the same layout in memory and on disk. output.write_device(short_factors_batch_.data_handle(), short_factors_size); output.write_device(long_code_.data_handle(), long_code_size); output.write_device(ex_factor_.data_handle(), ex_factor_size); diff --git a/cpp/src/util/serialize_validation.hpp b/cpp/src/util/serialize_validation.hpp index 0e8faea7c6..7fa49e8e2b 100644 --- a/cpp/src/util/serialize_validation.hpp +++ b/cpp/src/util/serialize_validation.hpp @@ -44,6 +44,14 @@ inline bool validate_serialized_dtype(const char* dtype_prefix, std::size_t dtyp { if (dtype_prefix == nullptr || dtype_prefix_size != 4) { return false; } + if constexpr (std::is_same_v) { + // Older cuVS writers used RAFT's legacy half descriptor instead of NumPy's "(); expected_dtype.resize(dtype_prefix_size, '\0'); diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 2eed6a3853..db94f70637 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -473,7 +473,7 @@ ConfigureTest( ConfigureTest( NAME UTIL_TEST - PATH util/file_io_test.cpp util/host_memory_test.cpp + PATH util/file_io_test.cpp util/host_memory_test.cpp util/serialize_validation_test.cpp GPUS 1 PERCENT 100 ) diff --git a/cpp/tests/neighbors/ann_hnsw_ace.cuh b/cpp/tests/neighbors/ann_hnsw_ace.cuh index 90a3e46caf..1776abe11d 100644 --- a/cpp/tests/neighbors/ann_hnsw_ace.cuh +++ b/cpp/tests/neighbors/ann_hnsw_ace.cuh @@ -33,6 +33,18 @@ namespace cuvs::neighbors::hnsw { +enum class LayeredSource { + disk_ace, + regular, + inmem_ace, + device_standard_attached, + device_standard_explicit, + device_padded_attached, + device_padded_explicit, + host_standard, + host_padded +}; + struct AnnHnswAceInputs { int n_queries; int n_rows; @@ -444,7 +456,7 @@ class AnnHnswAceTest : public ::testing::TestWithParam { std::filesystem::remove_all(temp_dir); } - void testHnswAceLayeredBuildDeserializeSearch() + void testHnswAceLayeredBuildDeserializeSearch(LayeredSource source = LayeredSource::disk_ace) { size_t queries_size = ps.n_queries * ps.k; std::vector indexes_naive(queries_size); @@ -503,13 +515,13 @@ class AnnHnswAceTest : public ::testing::TestWithParam { 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; + if (source == LayeredSource::disk_ace || source == LayeredSource::inmem_ace) { + auto ace_params = graph_build_params::ace_params(); + ace_params.npartitions = ps.npartitions; + ace_params.use_disk = source == LayeredSource::disk_ace; + if (ace_params.use_disk) { ace_params.build_dir = temp_dir; } + hnsw_params.graph_build_params = ace_params; + } auto invalid_params = hnsw_params; invalid_params.hierarchy = hnsw::HnswHierarchy::CPU; @@ -517,8 +529,57 @@ class AnnHnswAceTest : public ::testing::TestWithParam { hnsw::build(handle_, invalid_params, raft::make_const_mdspan(database_host.view())), std::exception); - auto hnsw_index = - hnsw::build(handle_, hnsw_params, raft::make_const_mdspan(database_host.view())); + std::unique_ptr> hnsw_index; + auto host_vectors = raft::make_const_mdspan(database_host.view()); + if (source == LayeredSource::disk_ace || source == LayeredSource::inmem_ace || + source == LayeredSource::regular) { + hnsw_index = hnsw::build(handle_, hnsw_params, host_vectors); + } else { + cagra::index_params cagra_params; + cagra_params.metric = ps.metric; + cagra_params.graph_degree = 64; + cagra_params.intermediate_graph_degree = 128; + // Explicit-vector cases also cover CAGRA indexes without an attached dataset. + cagra_params.attach_dataset_on_build = source != LayeredSource::device_standard_explicit && + source != LayeredSource::device_padded_explicit; + auto device_vectors = + raft::make_device_matrix_view(database_dev.data(), ps.n_rows, ps.dim); + // Each CAGRA source goes out of scope before the artifact is read. + if (source == LayeredSource::device_standard_attached || + source == LayeredSource::device_standard_explicit) { + auto cagra_index = + cagra::build(handle_, cagra_params, make_device_standard_dataset_view(device_vectors)); + hnsw_index = source == LayeredSource::device_standard_attached + ? hnsw::from_cagra(handle_, hnsw_params, cagra_index) + : hnsw::from_cagra(handle_, hnsw_params, cagra_index, host_vectors); + } else if (source == LayeredSource::device_padded_attached || + source == LayeredSource::device_padded_explicit) { + cuvs::neighbors::test::padded_device_matrix_for_cagra padded(handle_, + device_vectors); + auto cagra_index = cagra::build(handle_, cagra_params, padded.view); + hnsw_index = source == LayeredSource::device_padded_attached + ? hnsw::from_cagra(handle_, hnsw_params, cagra_index) + : hnsw::from_cagra(handle_, hnsw_params, cagra_index, host_vectors); + } else if (source == LayeredSource::host_standard) { + auto cagra_index = + cagra::build(handle_, cagra_params, make_host_standard_dataset_view(host_vectors)); + EXPECT_THROW(hnsw::from_cagra(handle_, hnsw_params, cagra_index), std::exception); + auto wrong_rows = raft::make_host_matrix_view( + database_host.data_handle(), ps.n_rows - 1, ps.dim); + auto wrong_columns = raft::make_host_matrix_view( + database_host.data_handle(), ps.n_rows, ps.dim - 1); + EXPECT_THROW(hnsw::from_cagra(handle_, hnsw_params, cagra_index, wrong_rows), + std::exception); + EXPECT_THROW(hnsw::from_cagra(handle_, hnsw_params, cagra_index, wrong_columns), + std::exception); + hnsw_index = hnsw::from_cagra(handle_, hnsw_params, cagra_index, host_vectors); + } else { + auto padded = make_host_padded_dataset(handle_, database_host.view()); + auto cagra_index = cagra::build(handle_, cagra_params, padded->as_dataset_view()); + EXPECT_THROW(hnsw::from_cagra(handle_, hnsw_params, cagra_index), std::exception); + hnsw_index = hnsw::from_cagra(handle_, hnsw_params, cagra_index, host_vectors); + } + } ASSERT_NE(hnsw_index, nullptr); EXPECT_EQ(hnsw_index->hierarchy(), hnsw::HnswHierarchy::GPU); EXPECT_EQ(hnsw_index->output_format(), hnsw::HnswOutputFormat::GRAPH_ONLY); @@ -531,7 +592,7 @@ class AnnHnswAceTest : public ::testing::TestWithParam { for (const auto& entry : std::filesystem::directory_iterator(temp_dir)) { if (entry.path().extension() == ".cuvs") { ++cuvs_artifact_count; } } - EXPECT_EQ(cuvs_artifact_count, 1); + EXPECT_EQ(cuvs_artifact_count, source == LayeredSource::disk_ace ? 1 : 0); EXPECT_FALSE(std::filesystem::exists(std::filesystem::path(temp_dir) / "layered_hnsw")); auto indexes_hnsw_host = raft::make_host_matrix(ps.n_queries, ps.k); @@ -605,6 +666,17 @@ class AnnHnswAceTest : public ::testing::TestWithParam { hnsw::deserialize(handle_, artifact_path, float_dataset_file, &mixed_precision_index); ASSERT_NE(mixed_precision_index, nullptr); std::unique_ptr> mixed_precision_guard(mixed_precision_index); + const auto mixed_artifact = (std::filesystem::path(temp_dir) / "mixed_resaved.cuvs").string(); + hnsw::serialize(handle_, mixed_artifact, *mixed_precision_guard); + std::ifstream mixed_file(mixed_artifact, std::ios::binary); + cuvs::neighbors::hnsw::detail::layered_hnsw_file_header mixed_header{}; + mixed_file.read(reinterpret_cast(&mixed_header), sizeof(mixed_header)); + ASSERT_TRUE(mixed_file.good()); + EXPECT_EQ(mixed_header.construction_dtype, + static_cast(cuvs::neighbors::hnsw::detail::layered_hnsw_dtype::int8)); + hnsw::index* mixed_roundtrip = nullptr; + hnsw::deserialize(handle_, mixed_artifact, float_dataset_file, &mixed_roundtrip); + mixed_precision_guard.reset(mixed_roundtrip); hnsw::search(handle_, search_params, @@ -639,10 +711,34 @@ class AnnHnswAceTest : public ::testing::TestWithParam { } EXPECT_EQ(copied_file_count, 1); + const auto original_artifact_size = std::filesystem::file_size(copied_artifact); + hnsw_index.reset(); + if (source == LayeredSource::disk_ace) { + // The caller owns persistent disk ACE artifacts. + EXPECT_TRUE(std::filesystem::exists(artifact_path)); + std::filesystem::remove(artifact_path); + std::filesystem::remove(std::filesystem::path(temp_dir) / "cagra_graph.npy"); + std::filesystem::remove(std::filesystem::path(temp_dir) / "reordered_dataset.npy"); + } else { + EXPECT_FALSE(std::filesystem::exists(artifact_path)); + EXPECT_FALSE(std::filesystem::exists(std::filesystem::path(artifact_path).parent_path())); + } + + // A loaded graph remains serializable after its source artifact has been removed. + // Reuse the destination to exercise serialize's overwrite behavior as well. + hnsw::serialize(handle_, copied_artifact, *deserialized_guard); + EXPECT_EQ(std::filesystem::file_size(copied_artifact), original_artifact_size); + deserialized_guard.reset(); + hnsw::index* copied_index = nullptr; hnsw::deserialize(handle_, copied_artifact, dataset_file, &copied_index); ASSERT_NE(copied_index, nullptr); std::unique_ptr> copied_guard(copied_index); + hnsw::serialize(handle_, copied_artifact, *copied_guard); + hnsw::index* roundtrip_index = nullptr; + hnsw::deserialize(handle_, copied_artifact, dataset_file, &roundtrip_index); + copied_guard.reset(roundtrip_index); + EXPECT_EQ(copied_guard->output_format(), hnsw::HnswOutputFormat::GRAPH_ONLY); hnsw::search(handle_, search_params, 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..066fdedb16 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,68 @@ INSTANTIATE_TEST_CASE_P(AnnHnswInmemSpillTest, AnnHnswInmemSpillTest_float, ::testing::ValuesIn(hnsw_inmem_spill_inputs)); +// One nonaligned shape per scalar; avoid crossing the full ACE parameter matrix. +typedef AnnHnswAceTest AnnHnswLayeredSourcesTest_float; +TEST_P(AnnHnswLayeredSourcesTest_float, RegularBuild) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::regular); +} + +TEST_P(AnnHnswLayeredSourcesTest_float, InMemoryAceBuild) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::inmem_ace); +} + +TEST_P(AnnHnswLayeredSourcesTest_float, DeviceStandardAttached) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_standard_attached); +} + +TEST_P(AnnHnswLayeredSourcesTest_float, DeviceStandardExplicit) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_standard_explicit); +} + +TEST_P(AnnHnswLayeredSourcesTest_float, DevicePaddedAttached) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_padded_attached); +} + +TEST_P(AnnHnswLayeredSourcesTest_float, DevicePaddedExplicit) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_padded_explicit); +} + +TEST_P(AnnHnswLayeredSourcesTest_float, HostStandard) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::host_standard); +} + +TEST_P(AnnHnswLayeredSourcesTest_float, HostPadded) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::host_padded); +} + +INSTANTIATE_TEST_CASE_P( + AnnHnswLayeredSourcesTest, + AnnHnswLayeredSourcesTest_float, + ::testing::Values(AnnHnswAceInputs{ + 10, 2000, 17, 10, 2, 100, false, cuvs::distance::DistanceType::L2Expanded, 0.9})); + +// Keep a small input regression alongside the larger conversion cases. +typedef AnnHnswAceTest AnnHnswLayeredTinyTest_float; +TEST_P(AnnHnswLayeredTinyTest_float, RegularBuild) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::regular); +} +TEST_P(AnnHnswLayeredTinyTest_float, InMemoryAceBuild) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::inmem_ace); +} +INSTANTIATE_TEST_CASE_P( + AnnHnswLayeredTinyTest, + AnnHnswLayeredTinyTest_float, + ::testing::Values(AnnHnswAceInputs{ + 10, 256, 17, 10, 2, 100, false, cuvs::distance::DistanceType::L2Expanded, 0.9})); + } // 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..702611162b 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,52 @@ INSTANTIATE_TEST_CASE_P(AnnHnswInmemSpillTest, AnnHnswInmemSpillTest_half, ::testing::ValuesIn(hnsw_inmem_spill_inputs)); +// One nonaligned shape per scalar; avoid crossing the full ACE parameter matrix. +typedef AnnHnswAceTest AnnHnswLayeredSourcesTest_half; +TEST_P(AnnHnswLayeredSourcesTest_half, RegularBuild) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::regular); +} + +TEST_P(AnnHnswLayeredSourcesTest_half, InMemoryAceBuild) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::inmem_ace); +} + +TEST_P(AnnHnswLayeredSourcesTest_half, DeviceStandardAttached) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_standard_attached); +} + +TEST_P(AnnHnswLayeredSourcesTest_half, DeviceStandardExplicit) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_standard_explicit); +} + +TEST_P(AnnHnswLayeredSourcesTest_half, DevicePaddedAttached) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_padded_attached); +} + +TEST_P(AnnHnswLayeredSourcesTest_half, DevicePaddedExplicit) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_padded_explicit); +} + +TEST_P(AnnHnswLayeredSourcesTest_half, HostStandard) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::host_standard); +} + +TEST_P(AnnHnswLayeredSourcesTest_half, HostPadded) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::host_padded); +} + +INSTANTIATE_TEST_CASE_P( + AnnHnswLayeredSourcesTest, + AnnHnswLayeredSourcesTest_half, + ::testing::Values(AnnHnswAceInputs{ + 10, 2000, 17, 10, 2, 100, false, cuvs::distance::DistanceType::L2Expanded, 0.9})); + } // 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..73b0f73dc9 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,52 @@ INSTANTIATE_TEST_CASE_P(AnnHnswInmemSpillTest, AnnHnswInmemSpillTest_int8_t, ::testing::ValuesIn(hnsw_inmem_spill_inputs)); +// One nonaligned shape per scalar; avoid crossing the full ACE parameter matrix. +typedef AnnHnswAceTest AnnHnswLayeredSourcesTest_int8_t; +TEST_P(AnnHnswLayeredSourcesTest_int8_t, RegularBuild) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::regular); +} + +TEST_P(AnnHnswLayeredSourcesTest_int8_t, InMemoryAceBuild) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::inmem_ace); +} + +TEST_P(AnnHnswLayeredSourcesTest_int8_t, DeviceStandardAttached) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_standard_attached); +} + +TEST_P(AnnHnswLayeredSourcesTest_int8_t, DeviceStandardExplicit) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_standard_explicit); +} + +TEST_P(AnnHnswLayeredSourcesTest_int8_t, DevicePaddedAttached) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_padded_attached); +} + +TEST_P(AnnHnswLayeredSourcesTest_int8_t, DevicePaddedExplicit) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_padded_explicit); +} + +TEST_P(AnnHnswLayeredSourcesTest_int8_t, HostStandard) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::host_standard); +} + +TEST_P(AnnHnswLayeredSourcesTest_int8_t, HostPadded) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::host_padded); +} + +INSTANTIATE_TEST_CASE_P( + AnnHnswLayeredSourcesTest, + AnnHnswLayeredSourcesTest_int8_t, + ::testing::Values(AnnHnswAceInputs{ + 10, 2000, 17, 10, 2, 100, false, cuvs::distance::DistanceType::L2Expanded, 0.9})); + } // 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..de54bd1dc8 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,52 @@ INSTANTIATE_TEST_CASE_P(AnnHnswInmemSpillTest, AnnHnswInmemSpillTest_uint8_t, ::testing::ValuesIn(hnsw_inmem_spill_inputs)); +// One nonaligned shape per scalar; avoid crossing the full ACE parameter matrix. +typedef AnnHnswAceTest AnnHnswLayeredSourcesTest_uint8_t; +TEST_P(AnnHnswLayeredSourcesTest_uint8_t, RegularBuild) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::regular); +} + +TEST_P(AnnHnswLayeredSourcesTest_uint8_t, InMemoryAceBuild) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::inmem_ace); +} + +TEST_P(AnnHnswLayeredSourcesTest_uint8_t, DeviceStandardAttached) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_standard_attached); +} + +TEST_P(AnnHnswLayeredSourcesTest_uint8_t, DeviceStandardExplicit) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_standard_explicit); +} + +TEST_P(AnnHnswLayeredSourcesTest_uint8_t, DevicePaddedAttached) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_padded_attached); +} + +TEST_P(AnnHnswLayeredSourcesTest_uint8_t, DevicePaddedExplicit) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::device_padded_explicit); +} + +TEST_P(AnnHnswLayeredSourcesTest_uint8_t, HostStandard) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::host_standard); +} + +TEST_P(AnnHnswLayeredSourcesTest_uint8_t, HostPadded) +{ + this->testHnswAceLayeredBuildDeserializeSearch(LayeredSource::host_padded); +} + +INSTANTIATE_TEST_CASE_P( + AnnHnswLayeredSourcesTest, + AnnHnswLayeredSourcesTest_uint8_t, + ::testing::Values(AnnHnswAceInputs{ + 10, 2000, 17, 10, 2, 100, false, cuvs::distance::DistanceType::L2Expanded, 0.9})); + } // namespace cuvs::neighbors::hnsw diff --git a/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu b/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu index 42bf84e6a3..9becc514ce 100644 --- a/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu +++ b/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu @@ -3,10 +3,101 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "../../../src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cuh" #include "../ann_ivf_rabitq.cuh" +#include + namespace cuvs::neighbors::ivf_rabitq { +// Check the historical disk layout independently of the save/load round trip: +// matching mistakes in both directions must not make this test pass. +TEST(IvfRabitqSerialization, PreservesLegacyVectorMajorCodes) +{ + raft::resources handle; + auto stream = raft::resource::get_cuda_stream(handle); + constexpr int64_t n_rows = 256; + constexpr int64_t dim = 128; + auto dataset = raft::make_device_matrix(handle, n_rows, dim); + raft::random::RngState rng(1234ULL); + raft::random::uniform(handle, rng, dataset.data_handle(), n_rows * dim, 0.1f, 2.0f); + index_params params; + params.n_lists = 4; + auto built = build(handle, params, raft::make_const_mdspan(dataset.view())); + auto& impl = built.rabitq_index(); + const size_t words = impl.quantizer().short_code_length(); + ASSERT_GT(words, 1); + std::vector clusters(params.n_lists); + raft::copy(clusters.data(), impl.get_cluster_meta().data_handle(), clusters.size(), stream); + raft::resource::sync_stream(handle); + + // Distinct words expose transposition mistakes even when quantized inputs happen + // to have identical bit patterns. Include every cluster's independent offset. + std::vector memory_codes(n_rows * words); + std::vector legacy_codes(n_rows * words); + bool has_multiple_rows = false; + for (auto const& cluster : clusters) { + has_multiple_rows |= cluster.num > 1; + const size_t offset = cluster.start_index * words; + for (size_t row = 0; row < cluster.num; ++row) { + for (size_t word = 0; word < words; ++word) { + const uint32_t value = static_cast(offset + row * words + word + 1); + legacy_codes[offset + row * words + word] = value; + memory_codes[offset + word * cluster.num + row] = value; + } + } + } + ASSERT_TRUE(has_multiple_rows); + raft::copy(impl.get_short_data_device(), memory_codes.data(), memory_codes.size(), stream); + raft::resource::sync_stream(handle); + tmp_index_file file; + serialize(handle, file.filename, built); + + // Short codes precede factors, long codes, extra factors, and IDs at EOF. + const size_t trailing_bytes = n_rows * (3 * sizeof(float) + impl.quantizer().long_code_length() + + 2 * sizeof(float) + sizeof(uint32_t)); + std::ifstream input(file.filename, std::ios::binary); + input.seekg(-static_cast(trailing_bytes + legacy_codes.size() * sizeof(uint32_t)), + std::ios::end); + std::vector saved_codes(legacy_codes.size()); + input.read(reinterpret_cast(saved_codes.data()), saved_codes.size() * sizeof(uint32_t)); + ASSERT_TRUE(input.good()); + EXPECT_EQ(saved_codes, legacy_codes); + input.close(); + + // Supply known legacy payload bytes independently of the writer under test. + std::fstream fixture(file.filename, std::ios::binary | std::ios::in | std::ios::out); + fixture.seekp( + -static_cast(trailing_bytes + legacy_codes.size() * sizeof(uint32_t)), + std::ios::end); + fixture.write(reinterpret_cast(legacy_codes.data()), + legacy_codes.size() * sizeof(uint32_t)); + ASSERT_TRUE(fixture.good()); + fixture.close(); + + index loaded(handle); + deserialize(handle, file.filename, &loaded); + std::vector loaded_codes(memory_codes.size()); + raft::copy(loaded_codes.data(), + loaded.rabitq_index().get_short_data_device(), + loaded_codes.size(), + stream); + raft::resource::sync_stream(handle); + EXPECT_EQ(loaded_codes, memory_codes); + + // Re-saving must also preserve the historical format rather than transposing twice. + tmp_index_file resaved; + serialize(handle, resaved.filename, loaded); + std::ifstream resaved_input(resaved.filename, std::ios::binary); + resaved_input.seekg( + -static_cast(trailing_bytes + legacy_codes.size() * sizeof(uint32_t)), + std::ios::end); + resaved_input.read(reinterpret_cast(saved_codes.data()), + saved_codes.size() * sizeof(uint32_t)); + ASSERT_TRUE(resaved_input.good()); + EXPECT_EQ(saved_codes, legacy_codes); +} + using f32_f32_i64 = ivf_rabitq_test; TEST_BUILD_SEARCH(f32_f32_i64) diff --git a/cpp/tests/util/serialize_validation_test.cpp b/cpp/tests/util/serialize_validation_test.cpp new file mode 100644 index 0000000000..644eb10106 --- /dev/null +++ b/cpp/tests/util/serialize_validation_test.cpp @@ -0,0 +1,44 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "../../src/util/serialize_validation.hpp" + +#include + +namespace cuvs::util { + +TEST(SerializeValidation, AcceptsLegacyAndStandardHalfPrefixes) +{ + // Literal bytes from the old and current writers, independent of the dtype helper. + constexpr char legacy[] = {'<', 'e', '2', '\0'}; + constexpr char standard[] = {'<', 'f', '2', '\0'}; + EXPECT_TRUE(validate_serialized_dtype(legacy, sizeof(legacy))); + EXPECT_TRUE(validate_serialized_dtype(standard, sizeof(standard))); + EXPECT_EQ(detail::numpy_dtype_string(), "(nullptr, 4)); + EXPECT_FALSE(validate_serialized_dtype("("("("(">e2", 4)); + EXPECT_FALSE(validate_serialized_dtype(">f2", 4)); + EXPECT_FALSE(validate_serialized_dtype("("("("("("(" std::string write_local_dataset(raft::host_matrix_view dataset, const std::string& path) @@ -90,7 +92,8 @@ auto quantize_dataset(raft::device_resources const& dev_resources, return dataset_i8; } -auto make_hnsw_ace_params(const std::string& build_dir) -> cuvs::neighbors::hnsw::index_params +auto make_hnsw_params(const std::string& backend, const std::string& build_dir) + -> cuvs::neighbors::hnsw::index_params { using namespace cuvs::neighbors; @@ -101,11 +104,13 @@ auto make_hnsw_ace_params(const std::string& build_dir) -> cuvs::neighbors::hnsw hnsw_params.M = 32; hnsw_params.ef_construction = 120; - auto ace_params = hnsw::graph_build_params::ace_params(); - ace_params.npartitions = 4; - ace_params.use_disk = true; - ace_params.build_dir = build_dir; - hnsw_params.graph_build_params = ace_params; + if (backend != "regular") { + auto ace_params = hnsw::graph_build_params::ace_params(); + ace_params.npartitions = 4; + ace_params.use_disk = backend == "disk-ace"; + ace_params.build_dir = build_dir; + hnsw_params.graph_build_params = ace_params; + } return hnsw_params; } @@ -113,16 +118,15 @@ auto make_hnsw_ace_params(const std::string& build_dir) -> cuvs::neighbors::hnsw template auto hnsw_build(raft::device_resources const& dev_resources, const cuvs::neighbors::hnsw::index_params& hnsw_params, - raft::host_matrix_view dataset) -> std::string + raft::host_matrix_view dataset, + const std::string& artifact_path) -> std::string { using namespace cuvs::neighbors; - auto hnsw_index = hnsw::build(dev_resources, hnsw_params, dataset); - const auto artifact_path = hnsw_index->file_path(); - if (artifact_path.empty()) { - throw std::runtime_error("Expected layered HNSW build to return an artifact path."); - } - std::cout << " hnsw_build: layered artifact written to " << artifact_path << std::endl; + auto hnsw_index = hnsw::build(dev_resources, hnsw_params, dataset); + // In-memory backends own a temporary artifact. Persist it before the handle is destroyed. + hnsw::serialize(dev_resources, artifact_path, *hnsw_index); + std::cout << " hnsw_build: layered artifact saved to " << artifact_path << std::endl; return artifact_path; } @@ -186,8 +190,15 @@ void hnsw_search(raft::device_resources const& dev_resources, } // namespace -int main() +int main(int argc, char** argv) { + const std::string backend = argc > 1 ? argv[1] : "regular"; + const std::string build_dir = argc > 2 ? argv[2] : "/tmp/hnsw_ace_layered"; + if (argc > 3 || (backend != "regular" && backend != "ace" && backend != "disk-ace")) { + std::cerr << "Usage: " << argv[0] << " [regular|ace|disk-ace] [output-directory]" << std::endl; + return 1; + } + std::cout << "Graph build backend: " << backend << std::endl; raft::device_resources dev_resources; rmm::mr::pool_memory_resource pool_mr(rmm::mr::get_current_device_resource_ref(), @@ -224,18 +235,18 @@ int main() auto queries_host_view = raft::make_host_matrix_view( queries_host.data_handle(), n_queries, n_dim); - std::filesystem::create_directories(kBuildDir); + std::filesystem::create_directories(build_dir); #if HNSW_ACE_LAYERED_USE_QUANTIZATION auto dataset_i8 = quantize_dataset(dev_resources, dataset_host_view); auto dataset_i8_view = raft::make_host_matrix_view( dataset_i8.data_handle(), n_samples, n_dim); - auto dataset_path = - write_local_dataset(dataset_host_view, std::string{kBuildDir} + "/dataset.npy"); - auto hnsw_params = make_hnsw_ace_params(kBuildDir); + auto dataset_path = write_local_dataset(dataset_host_view, build_dir + "/dataset.npy"); + auto hnsw_params = make_hnsw_params(backend, build_dir); - std::cout << "[stage 2] Build graph-only HNSW artifact from int8 data with ACE" << std::endl; - auto artifact_path = hnsw_build(dev_resources, hnsw_params, dataset_i8_view); + std::cout << "[stage 2] Build graph-only HNSW artifact from int8 data" << std::endl; + auto artifact_path = hnsw_build( + dev_resources, hnsw_params, dataset_i8_view, build_dir + "/saved_graph.cuvs"); std::cout << "[stage 3] Attach original float dataset" << std::endl; auto hnsw_index = hnsw_deserialize(dev_resources, artifact_path, dataset_path); @@ -243,12 +254,12 @@ int main() std::cout << "[stage 4] Search float HNSW index" << std::endl; hnsw_search(dev_resources, *hnsw_index, queries_host_view); #else - auto dataset_path = - write_local_dataset(dataset_host_view, std::string{kBuildDir} + "/dataset.npy"); - auto hnsw_params = make_hnsw_ace_params(kBuildDir); + auto dataset_path = write_local_dataset(dataset_host_view, build_dir + "/dataset.npy"); + auto hnsw_params = make_hnsw_params(backend, build_dir); - 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 2] Build layered HNSW index" << std::endl; + auto artifact_path = hnsw_build( + dev_resources, hnsw_params, dataset_host_view, build_dir + "/saved_graph.cuvs"); std::cout << "[stage 3] Deserialize layered HNSW index" << std::endl; auto hnsw_index = hnsw_deserialize(dev_resources, artifact_path, dataset_path); diff --git a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md index cdf32501d4..d468bd4468 100644 --- a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md +++ b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md @@ -36,7 +36,9 @@ enum class HnswHierarchy { Output artifact format for an HNSW index -`HNSWLIB` produces the standard hnswlib index format. `GRAPH_ONLY` stores the graph separately from vectors. The current implementation requires `HnswHierarchy::GPU` and disk-backed ACE. Load a graph-only artifact with the two-filename `deserialize` overload, which reads the vectors from a separate local dataset. +`HNSWLIB` produces the standard hnswlib index format. `GRAPH_ONLY` stores the graph separately from vectors and requires `HnswHierarchy::GPU`. Regular CAGRA, in-memory ACE, and disk-backed ACE all support this format. Load a graph-only artifact with the two-filename `deserialize` overload and a separate dataset in original row order before searching. + +In-memory builds return a file-backed handle that owns a temporary artifact. Call `serialize` before destroying that handle to preserve the artifact. Disk-backed ACE artifacts remain in the caller's build directory after the handle is destroyed. ```cpp enum class HnswOutputFormat { @@ -71,7 +73,7 @@ struct index_params : cuvs::neighbors::index_params { | Name | Type | Description | | --- | --- | --- | | `hierarchy` | [`HnswHierarchy`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-hnswhierarchy) | Hierarchy build type for HNSW index when converting from CAGRA index | -| `output_format` | [`HnswOutputFormat`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-hnswoutputformat) | Output artifact format. Graph-only output currently requires a GPU hierarchy and disk-backed ACE. | +| `output_format` | [`HnswOutputFormat`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-hnswoutputformat) | Output artifact format. Graph-only output requires a GPU hierarchy. | | `ef_construction` | `int` | Size of the candidate list during hierarchy construction when hierarchy is `CPU` | | `num_threads` | `int` | Number of host threads to use to construct hierarchy when hierarchy is `CPU` or `GPU`. When the value is 0, the number of threads is automatically determined to the maximum number of threads available.
NOTE: When hierarchy is `GPU`, while the majority of the work is done on the GPU, initialization of the HNSW index itself and some other work is parallelized with the help of CPU threads. | | `M` | `size_t` | HNSW M parameter: number of bi-directional links per node (used when building with ACE). | @@ -185,6 +187,8 @@ Get file path for disk-backed index virtual std::string file_path() const; ``` +Graph-only artifacts from in-memory builds are temporary and remain valid while this handle is alive. Use `serialize` to save a persistent copy. Disk-backed ACE artifacts remain in the caller's build directory. + **Returns** `virtual std::string` @@ -226,6 +230,8 @@ The resulting graph is compatible for HNSW search, but is not an exact equivalen The HNSW index construction parameters `M` and `ef_construction` are the main parameters to control the graph degree and graph quality. We have additional options that can be used to fine tune graph building on the GPU (see `cuvs::neighbors::cagra::index_params`). In case the index does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is important to set `ace_params.build_dir` to a fast disk with sufficient storage size. +`GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. The returned handle references a graph artifact. Save it with `serialize`, then load it with the two-filename `deserialize` overload and an original-order dataset before searching. Temporary artifacts from in-memory builds are removed when the handle is destroyed. + NOTE: This function requires CUDA headers to be available at compile time. Usage example: @@ -257,6 +263,8 @@ The resulting graph is compatible for HNSW search, but is not an exact equivalen The HNSW index construction parameters `M` and `ef_construction` are the main parameters to control the graph degree and graph quality. We have additional options that can be used to fine tune graph building on the GPU (see `cuvs::neighbors::cagra::index_params`). In case the index does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is important to set `ace_params.build_dir` to a fast disk with sufficient storage size. +`GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. The returned handle references a graph artifact. Save it with `serialize`, then load it with the two-filename `deserialize` overload and an original-order dataset before searching. Temporary artifacts from in-memory builds are removed when the handle is destroyed. + NOTE: This function requires CUDA headers to be available at compile time. Usage example: @@ -288,6 +296,8 @@ The resulting graph is compatible for HNSW search, but is not an exact equivalen The HNSW index construction parameters `M` and `ef_construction` are the main parameters to control the graph degree and graph quality. We have additional options that can be used to fine tune graph building on the GPU (see `cuvs::neighbors::cagra::index_params`). In case the index does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is important to set `ace_params.build_dir` to a fast disk with sufficient storage size. +`GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. The returned handle references a graph artifact. Save it with `serialize`, then load it with the two-filename `deserialize` overload and an original-order dataset before searching. Temporary artifacts from in-memory builds are removed when the handle is destroyed. + NOTE: This function requires CUDA headers to be available at compile time. Usage example: @@ -319,6 +329,8 @@ The resulting graph is compatible for HNSW search, but is not an exact equivalen The HNSW index construction parameters `M` and `ef_construction` are the main parameters to control the graph degree and graph quality. We have additional options that can be used to fine tune graph building on the GPU (see `cuvs::neighbors::cagra::index_params`). In case the index does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is important to set `ace_params.build_dir` to a fast disk with sufficient storage size. +`GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. The returned handle references a graph artifact. Save it with `serialize`, then load it with the two-filename `deserialize` overload and an original-order dataset before searching. Temporary artifacts from in-memory builds are removed when the handle is destroyed. + NOTE: This function requires CUDA headers to be available at compile time. Usage example: @@ -344,7 +356,7 @@ Construct an hnswlib index from a CAGRA index NOTE: When `hnsw::index_params.hie 1. `NONE`: This method uses the filesystem to write the CAGRA index in `/tmp/<random_number>.bin` before reading it as an hnswlib index, then deleting the temporary file. The returned index is immutable and can only be searched by the hnswlib wrapper in cuVS, as the format is not compatible with the original hnswlib. 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The serialized index is also compatible with the original hnswlib library. -3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset. +3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset in original row order. In-memory conversion returns a handle owning a temporary artifact. Call `serialize` before destroying the handle to preserve it. ```cpp std::unique_ptr> from_cagra( @@ -355,6 +367,8 @@ std::optional> dat std::nullopt); ``` +`GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. + Usage example: **Parameters** @@ -364,7 +378,7 @@ Usage example: | `res` | in | `raft::resources const&` | raft resources | | `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | | `cagra_index` | in | `const cuvs::neighbors::cagra::device_padded_index&` | cagra index | -| `dataset` | in | `std::optional>` | optional dataset to avoid extra memory copy when hierarchy is `CPU`

Default: `std::nullopt`. | +| `dataset` | in | `std::optional>` | optional original-order host vectors. For a GPU hierarchy, including `GRAPH_ONLY`, provide these if the CAGRA index has no attached device dataset. Rows and logical dimensions must match the CAGRA index. Attached padded device datasets are supported.

Default: `std::nullopt`. | **Returns** @@ -376,7 +390,7 @@ Construct an hnswlib index from a CAGRA index NOTE: When `hnsw::index_params.hie 1. `NONE`: This method uses the filesystem to write the CAGRA index in `/tmp/<random_number>.bin` before reading it as an hnswlib index, then deleting the temporary file. The returned index is immutable and can only be searched by the hnswlib wrapper in cuVS, as the format is not compatible with the original hnswlib. 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The serialized index is also compatible with the original hnswlib library. -3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset. +3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset in original row order. In-memory conversion returns a handle owning a temporary artifact. Call `serialize` before destroying the handle to preserve it. ```cpp std::unique_ptr> from_cagra( @@ -387,6 +401,8 @@ std::optional> data std::nullopt); ``` +`GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. + Usage example: **Parameters** @@ -396,7 +412,7 @@ Usage example: | `res` | in | `raft::resources const&` | raft resources | | `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | | `cagra_index` | in | `const cuvs::neighbors::cagra::device_padded_index&` | cagra index | -| `dataset` | in | `std::optional>` | optional dataset to avoid extra memory copy when hierarchy is `CPU`

Default: `std::nullopt`. | +| `dataset` | in | `std::optional>` | optional original-order host vectors. For a GPU hierarchy, including `GRAPH_ONLY`, provide these if the CAGRA index has no attached device dataset. Rows and logical dimensions must match the CAGRA index. Attached padded device datasets are supported.

Default: `std::nullopt`. | **Returns** @@ -408,7 +424,7 @@ Construct an hnswlib index from a CAGRA index NOTE: When `hnsw::index_params.hie 1. `NONE`: This method uses the filesystem to write the CAGRA index in `/tmp/<random_number>.bin` before reading it as an hnswlib index, then deleting the temporary file. The returned index is immutable and can only be searched by the hnswlib wrapper in cuVS, as the format is not compatible with the original hnswlib. 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The serialized index is also compatible with the original hnswlib library. -3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset. +3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset in original row order. In-memory conversion returns a handle owning a temporary artifact. Call `serialize` before destroying the handle to preserve it. ```cpp std::unique_ptr> from_cagra( @@ -419,6 +435,8 @@ std::optional> d std::nullopt); ``` +`GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. + Usage example: **Parameters** @@ -428,7 +446,7 @@ Usage example: | `res` | in | `raft::resources const&` | raft resources | | `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | | `cagra_index` | in | `const cuvs::neighbors::cagra::device_padded_index&` | cagra index | -| `dataset` | in | `std::optional>` | optional dataset to avoid extra memory copy when hierarchy is `CPU`

Default: `std::nullopt`. | +| `dataset` | in | `std::optional>` | optional original-order host vectors. For a GPU hierarchy, including `GRAPH_ONLY`, provide these if the CAGRA index has no attached device dataset. Rows and logical dimensions must match the CAGRA index. Attached padded device datasets are supported.

Default: `std::nullopt`. | **Returns** @@ -440,7 +458,7 @@ Construct an hnswlib index from a CAGRA index NOTE: When `hnsw::index_params.hie 1. `NONE`: This method uses the filesystem to write the CAGRA index in `/tmp/<random_number>.bin` before reading it as an hnswlib index, then deleting the temporary file. The returned index is immutable and can only be searched by the hnswlib wrapper in cuVS, as the format is not compatible with the original hnswlib. 2. `CPU`: The returned index is mutable and can be extended with additional vectors. The serialized index is also compatible with the original hnswlib library. -3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset. +3. `GPU`: The hierarchy is constructed on the GPU. When `output_format` is `GRAPH_ONLY`, the GPU-built hierarchy is stored as graph links only. Reload it with the two-filename `deserialize` overload so vectors can be reconstructed from the local dataset in original row order. In-memory conversion returns a handle owning a temporary artifact. Call `serialize` before destroying the handle to preserve it. ```cpp std::unique_ptr> from_cagra( @@ -451,6 +469,8 @@ std::optional> da std::nullopt); ``` +`GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. + Usage example: **Parameters** @@ -460,7 +480,7 @@ Usage example: | `res` | in | `raft::resources const&` | raft resources | | `params` | in | [`const index_params&`](/api-reference/cpp-api-neighbors-hnsw#neighbors-hnsw-index-params) | hnsw index parameters | | `cagra_index` | in | `const cuvs::neighbors::cagra::device_padded_index&` | cagra index | -| `dataset` | in | `std::optional>` | optional dataset to avoid extra memory copy when hierarchy is `CPU`

Default: `std::nullopt`. | +| `dataset` | in | `std::optional>` | optional original-order host vectors. For a GPU hierarchy, including `GRAPH_ONLY`, provide these if the CAGRA index has no attached device dataset. Rows and logical dimensions must match the CAGRA index. Attached padded device datasets are supported.

Default: `std::nullopt`. | **Returns** @@ -479,7 +499,9 @@ std::optional> dat std::nullopt); ``` -When the index has an attached device dataset view, `dataset` may be omitted. Otherwise pass a host matrix with the vectors (same contract as `device_padded_index`). +When the index has an attached device dataset view, `dataset` may be omitted. Otherwise pass a host matrix with the vectors (same contract as `device_padded_index`). `GRAPH_ONLY` requires a GPU hierarchy and follows the same dataset and temporary-artifact contracts. + +`GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. **Parameters** @@ -496,7 +518,7 @@ When the index has an attached device dataset view, `dataset` may be omitted. Ot **Additional overload:** `neighbors::hnsw::from_cagra` -Construct an hnswlib index from a host-built CAGRA index. Requires `dataset` for in-memory indices — host builds do not store vectors in the index. +Construct an hnswlib index from a host-built CAGRA index. Requires `dataset` for in-memory indices — host builds do not store vectors in the index. This also applies to `GRAPH_ONLY`, which requires a GPU hierarchy. Supply vectors in original row order with matching row count and logical dimensions. Serialize the returned temporary graph artifact before destroying the handle. ```cpp std::unique_ptr> from_cagra( @@ -507,6 +529,8 @@ std::optional> dat std::nullopt); ``` +`GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. + **Parameters** | Name | Direction | Type | Description | @@ -522,7 +546,7 @@ std::nullopt); **Additional overload:** `neighbors::hnsw::from_cagra` -Construct an hnswlib index from a host-built CAGRA index (standard dataset layout). Requires `dataset` for in-memory indices — host builds do not store vectors in the index. +Construct an hnswlib index from a host-built CAGRA index (standard dataset layout). Requires `dataset` for in-memory indices — host builds do not store vectors in the index. This also applies to `GRAPH_ONLY`, which requires a GPU hierarchy. Supply vectors in original row order with matching row count and logical dimensions. Serialize the returned temporary graph artifact before destroying the handle. ```cpp std::unique_ptr> from_cagra( @@ -533,6 +557,8 @@ std::optional> dat std::nullopt); ``` +`GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. + **Parameters** | Name | Direction | Type | Description | @@ -802,7 +828,7 @@ Usage example: ### neighbors::hnsw::serialize -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. Serialize to a filename outside the handle's temporary directory to preserve the artifact beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer ownership. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -824,7 +850,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::serialize` -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. Serialize to a filename outside the handle's temporary directory to preserve the artifact beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer ownership. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -846,7 +872,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::serialize` -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. Serialize to a filename outside the handle's temporary directory to preserve the artifact beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer ownership. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -868,7 +894,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::serialize` -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. Serialize to a filename outside the handle's temporary directory to preserve the artifact beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer ownership. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); From 31bcbd5edd68b6612b67649fd35ced6373e96003 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Wed, 23 Sep 2026 14:47:26 +0200 Subject: [PATCH 02/20] Get the stream from the reference --- cpp/src/neighbors/hnsw_layered_gather.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/neighbors/hnsw_layered_gather.cu b/cpp/src/neighbors/hnsw_layered_gather.cu index 073da3304a..d49dc3ea51 100644 --- a/cpp/src/neighbors/hnsw_layered_gather.cu +++ b/cpp/src/neighbors/hnsw_layered_gather.cu @@ -36,7 +36,7 @@ void gather_layered_hnsw_vectors(raft::resources const& res, { if (rows == 0) { return; } RAFT_EXPECTS(row_bytes > 0 && source_stride_bytes >= row_bytes, "Invalid vector row stride"); - const auto stream = raft::resource::get_cuda_stream(res); + cudaStream_t stream = raft::resource::get_cuda_stream(res).get(); const auto batch_rows = std::min(rows, std::max(1, (64 * 1024 * 1024) / (row_bytes + sizeof(size_t)))); rmm::device_uvector row_ids(batch_rows, stream); From 270364d8fcecdd26d339836fbf2439e3d86bbafd Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Fri, 9 Oct 2026 11:34:47 +0200 Subject: [PATCH 03/20] Document disk-backed HNSW graph-only conversion artifact precondition --- cpp/include/cuvs/neighbors/hnsw.hpp | 28 ++++++++++++++++++++ cpp/src/neighbors/detail/hnsw.hpp | 2 ++ fern/pages/cpp_api/cpp-api-neighbors-hnsw.md | 14 ++++++++++ 3 files changed, 44 insertions(+) diff --git a/cpp/include/cuvs/neighbors/hnsw.hpp b/cpp/include/cuvs/neighbors/hnsw.hpp index a50ba7ef48..adc2592a76 100644 --- a/cpp/include/cuvs/neighbors/hnsw.hpp +++ b/cpp/include/cuvs/neighbors/hnsw.hpp @@ -519,6 +519,10 @@ std::unique_ptr> build( * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails + * if that file already exists. Before repeating conversion, move the previous artifact elsewhere + * or use a fresh graph directory. + * * @param[in] res raft resources * @param[in] params hnsw index parameters * @param[in] cagra_index cagra index @@ -565,6 +569,10 @@ std::unique_ptr> from_cagra( * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails + * if that file already exists. Before repeating conversion, move the previous artifact elsewhere + * or use a fresh graph directory. + * * @param[in] res raft resources * @param[in] params hnsw index parameters * @param[in] cagra_index cagra index @@ -611,6 +619,10 @@ std::unique_ptr> from_cagra( * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails + * if that file already exists. Before repeating conversion, move the previous artifact elsewhere + * or use a fresh graph directory. + * * @param[in] res raft resources * @param[in] params hnsw index parameters * @param[in] cagra_index cagra index @@ -657,6 +669,10 @@ std::unique_ptr> from_cagra( * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails + * if that file already exists. Before repeating conversion, move the previous artifact elsewhere + * or use a fresh graph directory. + * * @param[in] res raft resources * @param[in] params hnsw index parameters * @param[in] cagra_index cagra index @@ -694,6 +710,10 @@ std::unique_ptr> from_cagra( * * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. + * + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails + * if that file already exists. Before repeating conversion, move the previous artifact elsewhere + * or use a fresh graph directory. */ std::unique_ptr> from_cagra( raft::resources const& res, @@ -732,6 +752,10 @@ std::unique_ptr> from_cagra( * * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. + * + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails + * if that file already exists. Before repeating conversion, move the previous artifact elsewhere + * or use a fresh graph directory. */ std::unique_ptr> from_cagra( raft::resources const& res, @@ -770,6 +794,10 @@ std::unique_ptr> from_cagra( * * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. + * + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails + * if that file already exists. Before repeating conversion, move the previous artifact elsewhere + * or use a fresh graph directory. */ std::unique_ptr> from_cagra( raft::resources const& res, diff --git a/cpp/src/neighbors/detail/hnsw.hpp b/cpp/src/neighbors/detail/hnsw.hpp index a77b5c6362..fa4e5b368b 100644 --- a/cpp/src/neighbors/detail/hnsw.hpp +++ b/cpp/src/neighbors/detail/hnsw.hpp @@ -2389,6 +2389,8 @@ std::unique_ptr> from_cagra( const auto graph_path = cagra_index.graph_fd()->get_path(); RAFT_EXPECTS(!graph_path.empty(), "Unable to get graph file path"); RAFT_EXPECTS(dataset.has_value(), "Layered HNSW requires the original-order host dataset"); + // Exclusive publication preserves existing artifacts. Repeated conversion requires + // moving the previous artifact elsewhere or using a fresh graph directory. artifact_path = std::filesystem::path(graph_path).parent_path() / "hnsw_index.cuvs"; } else { directory = std::make_unique(); diff --git a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md index d468bd4468..8973eb15a3 100644 --- a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md +++ b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md @@ -369,6 +369,8 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. + Usage example: **Parameters** @@ -403,6 +405,8 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. + Usage example: **Parameters** @@ -437,6 +441,8 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. + Usage example: **Parameters** @@ -471,6 +477,8 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. + Usage example: **Parameters** @@ -503,6 +511,8 @@ When the index has an attached device dataset view, `dataset` may be omitted. Ot `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. + **Parameters** | Name | Direction | Type | Description | @@ -531,6 +541,8 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. + **Parameters** | Name | Direction | Type | Description | @@ -559,6 +571,8 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. + **Parameters** | Name | Direction | Type | Description | From cbf4dd9bb114e1ba1651772fdf2dbbdbae25b8e3 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Fri, 9 Oct 2026 11:48:13 +0200 Subject: [PATCH 04/20] Remove unrelated RaBitQ changes from HNSW graph-only proposal --- .../neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu | 55 ++--------- .../ann_ivf_rabitq/test_float_int64_t.cu | 91 ------------------- 2 files changed, 8 insertions(+), 138 deletions(-) diff --git a/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu b/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu index 76d1de87d8..e27d34557a 100644 --- a/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu +++ b/cpp/src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu @@ -191,27 +191,8 @@ void IVFGPU::load(const char* filename) // New change: host copy of ivf. AllocateHostMemory(); - // The persisted format is vector-major within each cluster. Search and quantization - // use word-major codes, so convert at the serialization boundary. - const size_t words_per_vector = DQ->short_code_length(); - std::vector disk_codes(max_cluster_length * words_per_vector); - size_t code_offset = 0; - for (auto cluster_size : cluster_sizes) { - const size_t cluster_words = cluster_size * words_per_vector; - read_exact(disk_codes.data(), cluster_words * sizeof(uint32_t)); - for (size_t word = 0; word < words_per_vector; ++word) { - for (size_t row = 0; row < cluster_size; ++row) { - short_data_host_(code_offset + word * cluster_size + row) = - disk_codes[row * words_per_vector + word]; - } - } - code_offset += cluster_words; - } - raft::copy(short_data_.data_handle(), - short_data_host_.data_handle(), - GetShortDataBytesSimple() / sizeof(uint32_t), - stream_); - raft::resource::sync_stream(handle_); + read_into_device_host( + short_data_.data_handle(), short_data_host_.data_handle(), GetShortDataBytesSimple()); read_into_device(short_factors_batch_.data_handle(), GetShortDataFactorBytesBatch()); // SoA layout - no copy on CPU read_into_device_host( @@ -341,37 +322,17 @@ void IVFGPU::save(const char* filename) const this->initializer->SaveCentroids(output, filename); // Compute sizes for device arrays. - size_t long_code_size = GetLongCodeBytes(); - size_t ex_factor_size = GetExFactorBytes(); - size_t ids_size = GetPIDsBytes(); + size_t short_data_size = GetShortDataBytesSimple(); + size_t long_code_size = GetLongCodeBytes(); + size_t ex_factor_size = GetExFactorBytes(); + size_t ids_size = GetPIDsBytes(); // for batch data size_t short_factors_size = GetShortDataFactorBytesBatch(); raft::resource::sync_stream(handle_); - // Preserve the legacy vector-major disk layout while keeping word-major codes - // in memory. Bound staging memory to the largest cluster. - const size_t words_per_vector = DQ->short_code_length(); - const size_t max_cluster_size = *std::max_element(cluster_sizes.begin(), cluster_sizes.end()); - std::vector memory_codes(max_cluster_size * words_per_vector); - std::vector disk_codes(max_cluster_size * words_per_vector); - for (auto const& cluster : h_cluster_meta) { - const size_t cluster_words = cluster.num * words_per_vector; - if (cluster_words == 0) { continue; } - raft::copy(memory_codes.data(), - short_data_.data_handle() + cluster.start_index * words_per_vector, - cluster_words, - stream_); - raft::resource::sync_stream(handle_); - for (size_t row = 0; row < cluster.num; ++row) { - for (size_t word = 0; word < words_per_vector; ++word) { - disk_codes[row * words_per_vector + word] = memory_codes[word * cluster.num + row]; - } - } - write_exact(disk_codes.data(), cluster_words * sizeof(uint32_t)); - } - - // The remaining arrays have the same layout in memory and on disk. + // Write raw arrays to file. + output.write_device(short_data_.data_handle(), short_data_size); output.write_device(short_factors_batch_.data_handle(), short_factors_size); output.write_device(long_code_.data_handle(), long_code_size); output.write_device(ex_factor_.data_handle(), ex_factor_size); diff --git a/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu b/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu index 9becc514ce..42bf84e6a3 100644 --- a/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu +++ b/cpp/tests/neighbors/ann_ivf_rabitq/test_float_int64_t.cu @@ -3,101 +3,10 @@ * SPDX-License-Identifier: Apache-2.0 */ -#include "../../../src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cuh" #include "../ann_ivf_rabitq.cuh" -#include - namespace cuvs::neighbors::ivf_rabitq { -// Check the historical disk layout independently of the save/load round trip: -// matching mistakes in both directions must not make this test pass. -TEST(IvfRabitqSerialization, PreservesLegacyVectorMajorCodes) -{ - raft::resources handle; - auto stream = raft::resource::get_cuda_stream(handle); - constexpr int64_t n_rows = 256; - constexpr int64_t dim = 128; - auto dataset = raft::make_device_matrix(handle, n_rows, dim); - raft::random::RngState rng(1234ULL); - raft::random::uniform(handle, rng, dataset.data_handle(), n_rows * dim, 0.1f, 2.0f); - index_params params; - params.n_lists = 4; - auto built = build(handle, params, raft::make_const_mdspan(dataset.view())); - auto& impl = built.rabitq_index(); - const size_t words = impl.quantizer().short_code_length(); - ASSERT_GT(words, 1); - std::vector clusters(params.n_lists); - raft::copy(clusters.data(), impl.get_cluster_meta().data_handle(), clusters.size(), stream); - raft::resource::sync_stream(handle); - - // Distinct words expose transposition mistakes even when quantized inputs happen - // to have identical bit patterns. Include every cluster's independent offset. - std::vector memory_codes(n_rows * words); - std::vector legacy_codes(n_rows * words); - bool has_multiple_rows = false; - for (auto const& cluster : clusters) { - has_multiple_rows |= cluster.num > 1; - const size_t offset = cluster.start_index * words; - for (size_t row = 0; row < cluster.num; ++row) { - for (size_t word = 0; word < words; ++word) { - const uint32_t value = static_cast(offset + row * words + word + 1); - legacy_codes[offset + row * words + word] = value; - memory_codes[offset + word * cluster.num + row] = value; - } - } - } - ASSERT_TRUE(has_multiple_rows); - raft::copy(impl.get_short_data_device(), memory_codes.data(), memory_codes.size(), stream); - raft::resource::sync_stream(handle); - tmp_index_file file; - serialize(handle, file.filename, built); - - // Short codes precede factors, long codes, extra factors, and IDs at EOF. - const size_t trailing_bytes = n_rows * (3 * sizeof(float) + impl.quantizer().long_code_length() + - 2 * sizeof(float) + sizeof(uint32_t)); - std::ifstream input(file.filename, std::ios::binary); - input.seekg(-static_cast(trailing_bytes + legacy_codes.size() * sizeof(uint32_t)), - std::ios::end); - std::vector saved_codes(legacy_codes.size()); - input.read(reinterpret_cast(saved_codes.data()), saved_codes.size() * sizeof(uint32_t)); - ASSERT_TRUE(input.good()); - EXPECT_EQ(saved_codes, legacy_codes); - input.close(); - - // Supply known legacy payload bytes independently of the writer under test. - std::fstream fixture(file.filename, std::ios::binary | std::ios::in | std::ios::out); - fixture.seekp( - -static_cast(trailing_bytes + legacy_codes.size() * sizeof(uint32_t)), - std::ios::end); - fixture.write(reinterpret_cast(legacy_codes.data()), - legacy_codes.size() * sizeof(uint32_t)); - ASSERT_TRUE(fixture.good()); - fixture.close(); - - index loaded(handle); - deserialize(handle, file.filename, &loaded); - std::vector loaded_codes(memory_codes.size()); - raft::copy(loaded_codes.data(), - loaded.rabitq_index().get_short_data_device(), - loaded_codes.size(), - stream); - raft::resource::sync_stream(handle); - EXPECT_EQ(loaded_codes, memory_codes); - - // Re-saving must also preserve the historical format rather than transposing twice. - tmp_index_file resaved; - serialize(handle, resaved.filename, loaded); - std::ifstream resaved_input(resaved.filename, std::ios::binary); - resaved_input.seekg( - -static_cast(trailing_bytes + legacy_codes.size() * sizeof(uint32_t)), - std::ios::end); - resaved_input.read(reinterpret_cast(saved_codes.data()), - saved_codes.size() * sizeof(uint32_t)); - ASSERT_TRUE(resaved_input.good()); - EXPECT_EQ(saved_codes, legacy_codes); -} - using f32_f32_i64 = ivf_rabitq_test; TEST_BUILD_SEARCH(f32_f32_i64) From 79d4dae8ec1ace032ccf456392b35244e06d07a1 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:15:26 +0200 Subject: [PATCH 05/20] Clarify GRAPH_ONLY build behavior and consolidate its usage example --- cpp/include/cuvs/neighbors/hnsw.hpp | 31 +++++--------------- fern/pages/cpp_api/cpp-api-neighbors-hnsw.md | 12 ++------ 2 files changed, 11 insertions(+), 32 deletions(-) diff --git a/cpp/include/cuvs/neighbors/hnsw.hpp b/cpp/include/cuvs/neighbors/hnsw.hpp index adc2592a76..4ced004937 100644 --- a/cpp/include/cuvs/neighbors/hnsw.hpp +++ b/cpp/include/cuvs/neighbors/hnsw.hpp @@ -49,9 +49,14 @@ enum class HnswHierarchy { NONE, CPU, GPU }; * disk-backed ACE all support this format. Load a graph-only artifact with the two-filename * `deserialize` overload and a separate dataset in original row order before searching. * - * In-memory builds return a file-backed handle that owns a temporary artifact. Call `serialize` - * before destroying that handle to preserve the artifact. Disk-backed ACE artifacts remain in - * the caller's build directory after the handle is destroyed. + * With `GRAPH_ONLY`, `build` and `from_cagra` construct the hierarchy and write the graph-only + * file before returning. The returned index refers to that file; it is not yet searchable. + * For an in-memory CAGRA source, the index owns a temporary directory that is deleted when the + * index is destroyed. Use `serialize` to save the graph outside that directory. Disk-backed ACE + * output remains in the caller's build directory after the index is destroyed. + * + * For a complete build, save, load, and search example, see + * [hnsw_ace_layered_example.cu](https://github.com/NVIDIA/cuvs/blob/main/examples/cpp/src/hnsw_ace_layered_example.cu). */ enum class HnswOutputFormat { HNSWLIB, @@ -252,11 +257,6 @@ struct extend_params { * does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is * important to set `ace_params.build_dir` to a fast disk with sufficient storage size. * - * `GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. - * The returned handle references a graph artifact. Save it with `serialize`, then load it with - * the two-filename `deserialize` overload and an original-order dataset before searching. - * Temporary artifacts from in-memory builds are removed when the handle is destroyed. - * * NOTE: This function requires CUDA headers to be available at compile time. * * @param[in] res raft resources @@ -315,11 +315,6 @@ std::unique_ptr> build( * does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is * important to set `ace_params.build_dir` to a fast disk with sufficient storage size. * - * `GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. - * The returned handle references a graph artifact. Save it with `serialize`, then load it with - * the two-filename `deserialize` overload and an original-order dataset before searching. - * Temporary artifacts from in-memory builds are removed when the handle is destroyed. - * * NOTE: This function requires CUDA headers to be available at compile time. * * @param[in] res raft resources @@ -378,11 +373,6 @@ std::unique_ptr> build( * does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is * important to set `ace_params.build_dir` to a fast disk with sufficient storage size. * - * `GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. - * The returned handle references a graph artifact. Save it with `serialize`, then load it with - * the two-filename `deserialize` overload and an original-order dataset before searching. - * Temporary artifacts from in-memory builds are removed when the handle is destroyed. - * * NOTE: This function requires CUDA headers to be available at compile time. * * @param[in] res raft resources @@ -441,11 +431,6 @@ std::unique_ptr> build( * does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is * important to set `ace_params.build_dir` to a fast disk with sufficient storage size. * - * `GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. - * The returned handle references a graph artifact. Save it with `serialize`, then load it with - * the two-filename `deserialize` overload and an original-order dataset before searching. - * Temporary artifacts from in-memory builds are removed when the handle is destroyed. - * * NOTE: This function requires CUDA headers to be available at compile time. * * @param[in] res raft resources diff --git a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md index 8973eb15a3..37e5194b60 100644 --- a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md +++ b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md @@ -38,7 +38,9 @@ Output artifact format for an HNSW index `HNSWLIB` produces the standard hnswlib index format. `GRAPH_ONLY` stores the graph separately from vectors and requires `HnswHierarchy::GPU`. Regular CAGRA, in-memory ACE, and disk-backed ACE all support this format. Load a graph-only artifact with the two-filename `deserialize` overload and a separate dataset in original row order before searching. -In-memory builds return a file-backed handle that owns a temporary artifact. Call `serialize` before destroying that handle to preserve the artifact. Disk-backed ACE artifacts remain in the caller's build directory after the handle is destroyed. +With `GRAPH_ONLY`, `build` and `from_cagra` construct the hierarchy and write the graph-only file before returning. The returned index refers to that file; it is not yet searchable. For an in-memory CAGRA source, the index owns a temporary directory that is deleted when the index is destroyed. Use `serialize` to save the graph outside that directory. Disk-backed ACE output remains in the caller's build directory after the index is destroyed. + +For a complete build, save, load, and search example, see [hnsw_ace_layered_example.cu](https://github.com/NVIDIA/cuvs/blob/main/examples/cpp/src/hnsw_ace_layered_example.cu). ```cpp enum class HnswOutputFormat { @@ -230,8 +232,6 @@ The resulting graph is compatible for HNSW search, but is not an exact equivalen The HNSW index construction parameters `M` and `ef_construction` are the main parameters to control the graph degree and graph quality. We have additional options that can be used to fine tune graph building on the GPU (see `cuvs::neighbors::cagra::index_params`). In case the index does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is important to set `ace_params.build_dir` to a fast disk with sufficient storage size. -`GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. The returned handle references a graph artifact. Save it with `serialize`, then load it with the two-filename `deserialize` overload and an original-order dataset before searching. Temporary artifacts from in-memory builds are removed when the handle is destroyed. - NOTE: This function requires CUDA headers to be available at compile time. Usage example: @@ -263,8 +263,6 @@ The resulting graph is compatible for HNSW search, but is not an exact equivalen The HNSW index construction parameters `M` and `ef_construction` are the main parameters to control the graph degree and graph quality. We have additional options that can be used to fine tune graph building on the GPU (see `cuvs::neighbors::cagra::index_params`). In case the index does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is important to set `ace_params.build_dir` to a fast disk with sufficient storage size. -`GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. The returned handle references a graph artifact. Save it with `serialize`, then load it with the two-filename `deserialize` overload and an original-order dataset before searching. Temporary artifacts from in-memory builds are removed when the handle is destroyed. - NOTE: This function requires CUDA headers to be available at compile time. Usage example: @@ -296,8 +294,6 @@ The resulting graph is compatible for HNSW search, but is not an exact equivalen The HNSW index construction parameters `M` and `ef_construction` are the main parameters to control the graph degree and graph quality. We have additional options that can be used to fine tune graph building on the GPU (see `cuvs::neighbors::cagra::index_params`). In case the index does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is important to set `ace_params.build_dir` to a fast disk with sufficient storage size. -`GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. The returned handle references a graph artifact. Save it with `serialize`, then load it with the two-filename `deserialize` overload and an original-order dataset before searching. Temporary artifacts from in-memory builds are removed when the handle is destroyed. - NOTE: This function requires CUDA headers to be available at compile time. Usage example: @@ -329,8 +325,6 @@ The resulting graph is compatible for HNSW search, but is not an exact equivalen The HNSW index construction parameters `M` and `ef_construction` are the main parameters to control the graph degree and graph quality. We have additional options that can be used to fine tune graph building on the GPU (see `cuvs::neighbors::cagra::index_params`). In case the index does not fit the host or GPU memory, we would use disk as temporary storage. In such cases it is important to set `ace_params.build_dir` to a fast disk with sufficient storage size. -`GRAPH_ONLY` supports regular CAGRA, in-memory ACE, and disk-backed ACE with a GPU hierarchy. The returned handle references a graph artifact. Save it with `serialize`, then load it with the two-filename `deserialize` overload and an original-order dataset before searching. Temporary artifacts from in-memory builds are removed when the handle is destroyed. - NOTE: This function requires CUDA headers to be available at compile time. Usage example: From 2336b0494fef905379521b251f9871da8dce0711 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:16:40 +0200 Subject: [PATCH 06/20] Explain how disk-backed HNSW conversion selects its output directory --- cpp/include/cuvs/neighbors/hnsw.hpp | 56 ++++++++++++-------- fern/pages/cpp_api/cpp-api-neighbors-hnsw.md | 14 ++--- 2 files changed, 42 insertions(+), 28 deletions(-) diff --git a/cpp/include/cuvs/neighbors/hnsw.hpp b/cpp/include/cuvs/neighbors/hnsw.hpp index 4ced004937..71675c89c0 100644 --- a/cpp/include/cuvs/neighbors/hnsw.hpp +++ b/cpp/include/cuvs/neighbors/hnsw.hpp @@ -504,9 +504,11 @@ std::unique_ptr> build( * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * - * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails - * if that file already exists. Before repeating conversion, move the previous artifact elsewhere - * or use a fresh graph directory. + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, + * in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` + * derives this path from the CAGRA graph file; it does not select a new build directory. + * Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the + * previous artifact elsewhere or use a fresh graph directory. * * @param[in] res raft resources * @param[in] params hnsw index parameters @@ -554,9 +556,11 @@ std::unique_ptr> from_cagra( * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * - * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails - * if that file already exists. Before repeating conversion, move the previous artifact elsewhere - * or use a fresh graph directory. + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, + * in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` + * derives this path from the CAGRA graph file; it does not select a new build directory. + * Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the + * previous artifact elsewhere or use a fresh graph directory. * * @param[in] res raft resources * @param[in] params hnsw index parameters @@ -604,9 +608,11 @@ std::unique_ptr> from_cagra( * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * - * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails - * if that file already exists. Before repeating conversion, move the previous artifact elsewhere - * or use a fresh graph directory. + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, + * in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` + * derives this path from the CAGRA graph file; it does not select a new build directory. + * Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the + * previous artifact elsewhere or use a fresh graph directory. * * @param[in] res raft resources * @param[in] params hnsw index parameters @@ -654,9 +660,11 @@ std::unique_ptr> from_cagra( * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * - * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails - * if that file already exists. Before repeating conversion, move the previous artifact elsewhere - * or use a fresh graph directory. + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, + * in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` + * derives this path from the CAGRA graph file; it does not select a new build directory. + * Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the + * previous artifact elsewhere or use a fresh graph directory. * * @param[in] res raft resources * @param[in] params hnsw index parameters @@ -696,9 +704,11 @@ std::unique_ptr> from_cagra( * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * - * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails - * if that file already exists. Before repeating conversion, move the previous artifact elsewhere - * or use a fresh graph directory. + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, + * in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` + * derives this path from the CAGRA graph file; it does not select a new build directory. + * Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the + * previous artifact elsewhere or use a fresh graph directory. */ std::unique_ptr> from_cagra( raft::resources const& res, @@ -738,9 +748,11 @@ std::unique_ptr> from_cagra( * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * - * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails - * if that file already exists. Before repeating conversion, move the previous artifact elsewhere - * or use a fresh graph directory. + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, + * in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` + * derives this path from the CAGRA graph file; it does not select a new build directory. + * Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the + * previous artifact elsewhere or use a fresh graph directory. */ std::unique_ptr> from_cagra( raft::resources const& res, @@ -780,9 +792,11 @@ std::unique_ptr> from_cagra( * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. * - * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails - * if that file already exists. Before repeating conversion, move the previous artifact elsewhere - * or use a fresh graph directory. + * Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, + * in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` + * derives this path from the CAGRA graph file; it does not select a new build directory. + * Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the + * previous artifact elsewhere or use a fresh graph directory. */ std::unique_ptr> from_cagra( raft::resources const& res, diff --git a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md index 37e5194b60..16e57502e1 100644 --- a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md +++ b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md @@ -363,7 +363,7 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. -Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` derives this path from the CAGRA graph file; it does not select a new build directory. Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. Usage example: @@ -399,7 +399,7 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. -Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` derives this path from the CAGRA graph file; it does not select a new build directory. Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. Usage example: @@ -435,7 +435,7 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. -Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` derives this path from the CAGRA graph file; it does not select a new build directory. Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. Usage example: @@ -471,7 +471,7 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. -Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` derives this path from the CAGRA graph file; it does not select a new build directory. Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. Usage example: @@ -505,7 +505,7 @@ When the index has an attached device dataset view, `dataset` may be omitted. Ot `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. -Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` derives this path from the CAGRA graph file; it does not select a new build directory. Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. **Parameters** @@ -535,7 +535,7 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. -Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` derives this path from the CAGRA graph file; it does not select a new build directory. Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. **Parameters** @@ -565,7 +565,7 @@ std::nullopt); `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. -Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the CAGRA graph and fails if that file already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. +Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existing CAGRA graph, in the directory selected by `ace_params.build_dir` when that graph was built. `from_cagra` derives this path from the CAGRA graph file; it does not select a new build directory. Conversion fails if `hnsw_index.cuvs` already exists. Before repeating conversion, move the previous artifact elsewhere or use a fresh graph directory. **Parameters** From 103c236c56ba90cf27bba7d14fe0642b9ccb4618 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:17:10 +0200 Subject: [PATCH 07/20] Distinguish temporary HNSW output from persistent ACE build directories --- cpp/include/cuvs/neighbors/hnsw.hpp | 32 ++++++++++++-------- fern/pages/cpp_api/cpp-api-neighbors-hnsw.md | 8 ++--- 2 files changed, 24 insertions(+), 16 deletions(-) diff --git a/cpp/include/cuvs/neighbors/hnsw.hpp b/cpp/include/cuvs/neighbors/hnsw.hpp index 71675c89c0..fce44d6bda 100644 --- a/cpp/include/cuvs/neighbors/hnsw.hpp +++ b/cpp/include/cuvs/neighbors/hnsw.hpp @@ -1198,9 +1198,11 @@ void search(raft::resources const& res, * library. * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. * Load it with the two-filename `deserialize` overload and a local dataset in original row order. - * Serialize to a filename outside the handle's temporary directory to preserve the artifact - * beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer - * ownership. + * For an in-memory CAGRA source, conversion creates an internal temporary directory that is + * deleted when the returned index is destroyed. Serialize outside that directory to keep the + * graph file. Serializing to the index's current `file_path()` does not transfer ownership. + * A caller-supplied `ace_params.build_dir` used for disk-backed ACE is not an owned temporary + * directory: destroying the index does not remove that directory or its graph file. * * @param[in] res raft resources * @param[in] filename path to the serialized HNSW output @@ -1232,9 +1234,11 @@ void serialize(raft::resources const& res, const std::string& filename, const in * library. * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. * Load it with the two-filename `deserialize` overload and a local dataset in original row order. - * Serialize to a filename outside the handle's temporary directory to preserve the artifact - * beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer - * ownership. + * For an in-memory CAGRA source, conversion creates an internal temporary directory that is + * deleted when the returned index is destroyed. Serialize outside that directory to keep the + * graph file. Serializing to the index's current `file_path()` does not transfer ownership. + * A caller-supplied `ace_params.build_dir` used for disk-backed ACE is not an owned temporary + * directory: destroying the index does not remove that directory or its graph file. * * @param[in] res raft resources * @param[in] filename path to the serialized HNSW output @@ -1266,9 +1270,11 @@ void serialize(raft::resources const& res, const std::string& filename, const in * library. * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. * Load it with the two-filename `deserialize` overload and a local dataset in original row order. - * Serialize to a filename outside the handle's temporary directory to preserve the artifact - * beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer - * ownership. + * For an in-memory CAGRA source, conversion creates an internal temporary directory that is + * deleted when the returned index is destroyed. Serialize outside that directory to keep the + * graph file. Serializing to the index's current `file_path()` does not transfer ownership. + * A caller-supplied `ace_params.build_dir` used for disk-backed ACE is not an owned temporary + * directory: destroying the index does not remove that directory or its graph file. * * @param[in] res raft resources * @param[in] filename path to the serialized HNSW output @@ -1300,9 +1306,11 @@ void serialize(raft::resources const& res, const std::string& filename, const in * library. * When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. * Load it with the two-filename `deserialize` overload and a local dataset in original row order. - * Serialize to a filename outside the handle's temporary directory to preserve the artifact - * beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer - * ownership. + * For an in-memory CAGRA source, conversion creates an internal temporary directory that is + * deleted when the returned index is destroyed. Serialize outside that directory to keep the + * graph file. Serializing to the index's current `file_path()` does not transfer ownership. + * A caller-supplied `ace_params.build_dir` used for disk-backed ACE is not an owned temporary + * directory: destroying the index does not remove that directory or its graph file. * * @param[in] res raft resources * @param[in] filename path to the serialized HNSW output diff --git a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md index 16e57502e1..471d08e5f0 100644 --- a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md +++ b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md @@ -836,7 +836,7 @@ Usage example: ### neighbors::hnsw::serialize -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. Serialize to a filename outside the handle's temporary directory to preserve the artifact beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer ownership. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. For an in-memory CAGRA source, conversion creates an internal temporary directory that is deleted when the returned index is destroyed. Serialize outside that directory to keep the graph file. Serializing to the index's current `file_path()` does not transfer ownership. A caller-supplied `ace_params.build_dir` used for disk-backed ACE is not an owned temporary directory: destroying the index does not remove that directory or its graph file. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -858,7 +858,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::serialize` -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. Serialize to a filename outside the handle's temporary directory to preserve the artifact beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer ownership. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. For an in-memory CAGRA source, conversion creates an internal temporary directory that is deleted when the returned index is destroyed. Serialize outside that directory to keep the graph file. Serializing to the index's current `file_path()` does not transfer ownership. A caller-supplied `ace_params.build_dir` used for disk-backed ACE is not an owned temporary directory: destroying the index does not remove that directory or its graph file. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -880,7 +880,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::serialize` -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. Serialize to a filename outside the handle's temporary directory to preserve the artifact beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer ownership. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. For an in-memory CAGRA source, conversion creates an internal temporary directory that is deleted when the returned index is destroyed. Serialize outside that directory to keep the graph file. Serializing to the index's current `file_path()` does not transfer ownership. A caller-supplied `ace_params.build_dir` used for disk-backed ACE is not an owned temporary directory: destroying the index does not remove that directory or its graph file. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); @@ -902,7 +902,7 @@ Usage example: **Additional overload:** `neighbors::hnsw::serialize` -Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. Serialize to a filename outside the handle's temporary directory to preserve the artifact beyond the handle's lifetime. Serializing to its current `file_path()` does not transfer ownership. +Serialize the HNSW index to file NOTE: When hierarchy is `NONE`, the saved hnswlib index is immutable and can only be read by the hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. However, when hierarchy is `CPU`, the saved hnswlib index is compatible with the original hnswlib library. When `output_format` is `GRAPH_ONLY`, the saved artifact stores the graph only. Load it with the two-filename `deserialize` overload and a local dataset in original row order. For an in-memory CAGRA source, conversion creates an internal temporary directory that is deleted when the returned index is destroyed. Serialize outside that directory to keep the graph file. Serializing to the index's current `file_path()` does not transfer ownership. A caller-supplied `ace_params.build_dir` used for disk-backed ACE is not an owned temporary directory: destroying the index does not remove that directory or its graph file. ```cpp void serialize(raft::resources const& res, const std::string& filename, const index& idx); From 128596d0e9ea5483e51a5863b50ffd7a02af9c59 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:17:44 +0200 Subject: [PATCH 08/20] Document HNSW gather residency sizes and transfer contract --- .../neighbors/detail/hnsw_layered_gather.hpp | 24 +++++++++++++++++-- 1 file changed, 22 insertions(+), 2 deletions(-) diff --git a/cpp/src/neighbors/detail/hnsw_layered_gather.hpp b/cpp/src/neighbors/detail/hnsw_layered_gather.hpp index b3e491f1aa..a07f1bab92 100644 --- a/cpp/src/neighbors/detail/hnsw_layered_gather.hpp +++ b/cpp/src/neighbors/detail/hnsw_layered_gather.hpp @@ -9,8 +9,28 @@ #include namespace cuvs::neighbors::hnsw::detail { -// Gather only selected device rows to a compact host matrix using bounded device staging. -// Row IDs must be valid source rows; host_output must hold rows * row_bytes bytes. +/** + * @brief Copy selected device vectors to host memory for HNSW upper-layer construction. + * + * This wrapper manages host/device transfers and bounded device staging. Row selection is a + * generic gather operation; no HNSW topology is interpreted here. Only the logical row bytes + * are copied, excluding any padding in the source. Copies use the resource stream and complete + * before this function returns. A request with zero rows performs no work. + * + * @param[in] res Resources supplying the CUDA stream and current device. + * @param[in] source Device-accessible source storage on the current device. Each selected row + * must have at least `row_bytes` readable bytes at `source + row_id * source_stride_bytes`. + * @param[in] source_stride_bytes Physical distance between source rows in bytes; must be at + * least `row_bytes` for a nonempty request. + * @param[in] row_bytes Logical bytes per vector; must be positive for a nonempty request. + * @param[in] host_row_ids Host-resident array of `rows` source-row indices. Every index must + * identify a valid source row; bounds are the caller's responsibility. Order and duplicates + * are preserved. + * @param[in] rows Number of vectors to gather, not the source row count. + * @param[out] host_output Host-resident writable storage of at least `rows * row_bytes` bytes. + * Output rows are contiguous, with no padding. Input and output storage must remain valid + * until the call returns. + */ void gather_layered_hnsw_vectors(raft::resources const& res, const void* source, size_t source_stride_bytes, From 303cda4019919982413041a6a52d6a6a8781345c Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:19:21 +0200 Subject: [PATCH 09/20] Reuse RAFT matrix gather for HNSW device vector staging --- cpp/src/neighbors/hnsw_layered_gather.cu | 40 +++++++++-------- cpp/tests/CMakeLists.txt | 2 +- .../ann_hnsw_ace/test_float_uint32_t.cu | 43 +++++++++++++++++++ 3 files changed, 65 insertions(+), 20 deletions(-) diff --git a/cpp/src/neighbors/hnsw_layered_gather.cu b/cpp/src/neighbors/hnsw_layered_gather.cu index d49dc3ea51..bfa4822e86 100644 --- a/cpp/src/neighbors/hnsw_layered_gather.cu +++ b/cpp/src/neighbors/hnsw_layered_gather.cu @@ -7,23 +7,25 @@ #include #include #include +#include #include #include +#include +#include namespace cuvs::neighbors::hnsw::detail { namespace { -__global__ void gather_rows(const uint8_t* source, - size_t stride, - size_t row_bytes, - const size_t* row_ids, - size_t bytes, - uint8_t* output) -{ - for (size_t i = blockIdx.x * size_t(blockDim.x) + threadIdx.x; i < bytes; - i += size_t(blockDim.x) * gridDim.x) { - output[i] = source[size_t(row_ids[i / row_bytes]) * stride + i % row_bytes]; +// Present padded rows as a compact byte matrix to RAFT's iterator-based gather. +struct strided_source_byte { + const uint8_t* source; + size_t stride; + size_t row_bytes; + + __host__ __device__ uint8_t operator()(size_t offset) const + { + return source[(offset / row_bytes) * stride + offset % row_bytes]; } -} +}; } // namespace void gather_layered_hnsw_vectors(raft::resources const& res, @@ -36,7 +38,12 @@ void gather_layered_hnsw_vectors(raft::resources const& res, { if (rows == 0) { return; } RAFT_EXPECTS(row_bytes > 0 && source_stride_bytes >= row_bytes, "Invalid vector row stride"); - cudaStream_t stream = raft::resource::get_cuda_stream(res).get(); + cudaStream_t stream = raft::resource::get_cuda_stream(res).get(); + const auto source_bytes = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), + strided_source_byte{static_cast(source), source_stride_bytes, row_bytes}); + // The source has at least enough rows to include every requested row ID. + const auto source_rows = *std::max_element(host_row_ids, host_row_ids + rows) + size_t{1}; const auto batch_rows = std::min(rows, std::max(1, (64 * 1024 * 1024) / (row_bytes + sizeof(size_t)))); rmm::device_uvector row_ids(batch_rows, stream); @@ -48,13 +55,8 @@ void gather_layered_hnsw_vectors(raft::resources const& res, count * sizeof(size_t), cudaMemcpyHostToDevice, stream)); - gather_rows<<(65535, (count * row_bytes + 255) / 256), 256, 0, stream>>>( - static_cast(source), - source_stride_bytes, - row_bytes, - row_ids.data(), - count * row_bytes, - values.data()); + raft::matrix::gather( + source_bytes, row_bytes, source_rows, row_ids.data(), count, values.data(), stream); RAFT_CUDA_TRY(cudaGetLastError()); RAFT_CUDA_TRY(cudaMemcpyAsync(static_cast(host_output) + start * row_bytes, values.data(), diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index a29d20a813..17d31739b4 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -351,7 +351,7 @@ if(BUILD_CAGRA_HNSWLIB) ConfigureTest( NAME NEIGHBORS_ANN_HNSW_ACE_FLOAT_UINT32_TEST - PATH neighbors/ann_hnsw_ace/test_float_uint32_t.cu + PATH neighbors/ann_hnsw_ace/test_float_uint32_t.cu ../src/neighbors/hnsw_layered_gather.cu GPUS 1 PERCENT 100 ) 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 066fdedb16..c9a826e6bc 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 @@ -3,10 +3,53 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include "../../../src/neighbors/detail/hnsw_layered_gather.hpp" #include "../ann_hnsw_ace.cuh" namespace cuvs::neighbors::hnsw { +TEST(HnswLayeredGather, PreservesSelectedBytesWithAndWithoutPadding) +{ + raft::resources handle; + auto stream = raft::resource::get_cuda_stream(handle); + const std::vector row_ids{4, 1, 4, 0}; + constexpr size_t source_rows = 5; + // Cover nonaligned vectors with one-, two-, and four-byte scalar representations. + for (const size_t row_bytes : {size_t{17}, size_t{34}, size_t{68}}) { + for (const size_t padding : {size_t{0}, size_t{12}}) { + const size_t stride = row_bytes + padding; + std::vector source(source_rows * stride, 0xff); + for (size_t row = 0; row < source_rows; ++row) { + for (size_t col = 0; col < row_bytes; ++col) { + source[row * stride + col] = static_cast(row * row_bytes + col); + } + } + rmm::device_uvector device_source(source.size(), stream); + raft::copy(device_source.data(), source.data(), source.size(), stream); + std::vector actual(row_ids.size() * row_bytes); + detail::gather_layered_hnsw_vectors(handle, + device_source.data(), + stride, + row_bytes, + row_ids.data(), + row_ids.size(), + actual.data()); + std::vector expected; + for (const auto row : row_ids) { + expected.insert( + expected.end(), source.begin() + row * stride, source.begin() + row * stride + row_bytes); + } + EXPECT_EQ(actual, expected) << "row_bytes=" << row_bytes << " padding=" << padding; + } + } +} + +TEST(HnswLayeredGather, EmptySelectionDoesNotAccessStorage) +{ + raft::resources handle; + EXPECT_NO_THROW(detail::gather_layered_hnsw_vectors(handle, nullptr, 0, 0, nullptr, 0, nullptr)); +} + TEST(CagraAceWorkspace, FailurePreservesCallerDirectory) { test_ace_workspace_failure_preserves_caller_directory(); From 016bf54491f78e57fa204d7999542bf577d70cc2 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:20:41 +0200 Subject: [PATCH 10/20] Explain the HNSW layered test source variants --- cpp/tests/neighbors/ann_hnsw_ace.cuh | 23 ++++++++++++++--------- 1 file changed, 14 insertions(+), 9 deletions(-) diff --git a/cpp/tests/neighbors/ann_hnsw_ace.cuh b/cpp/tests/neighbors/ann_hnsw_ace.cuh index 1776abe11d..c0921b92f8 100644 --- a/cpp/tests/neighbors/ann_hnsw_ace.cuh +++ b/cpp/tests/neighbors/ann_hnsw_ace.cuh @@ -33,16 +33,21 @@ namespace cuvs::neighbors::hnsw { +// Input construction paths for the GRAPH_ONLY build/load/search tests. +// Standard datasets use contiguous logical rows; padded datasets may have a larger physical +// row stride. Attached cases obtain vectors from the device CAGRA index. Explicit cases build +// without an attached dataset and pass original-order host vectors to from_cagra instead. +// Host-built CAGRA indexes always require those explicit host vectors during conversion. enum class LayeredSource { - disk_ace, - regular, - inmem_ace, - device_standard_attached, - device_standard_explicit, - device_padded_attached, - device_padded_explicit, - host_standard, - host_padded + disk_ace, // hnsw::build with explicit ACE parameters and disk mode enabled. + regular, // Default hnsw::build; small inputs exercise ordinary CAGRA. + inmem_ace, // hnsw::build with explicit ACE parameters and disk mode disabled. + device_standard_attached, // Device-built CAGRA, standard rows, attached device vectors. + device_standard_explicit, // Device-built CAGRA, standard rows, explicit host vectors. + device_padded_attached, // Device-built CAGRA, padded rows, attached device vectors. + device_padded_explicit, // Device-built CAGRA, padded rows, explicit host vectors. + host_standard, // Host-built CAGRA, standard input rows, explicit host vectors. + host_padded // Host-built CAGRA, padded input rows, explicit host vectors. }; struct AnnHnswAceInputs { From d862e5e4ff47fcb703721c01de37d76b71ed6e8d Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:20:57 +0200 Subject: [PATCH 11/20] Describe graph build method choices in the HNSW example --- examples/cpp/src/hnsw_ace_layered_example.cu | 26 +++++++++++--------- 1 file changed, 15 insertions(+), 11 deletions(-) diff --git a/examples/cpp/src/hnsw_ace_layered_example.cu b/examples/cpp/src/hnsw_ace_layered_example.cu index eb96b62346..dd55763183 100644 --- a/examples/cpp/src/hnsw_ace_layered_example.cu +++ b/examples/cpp/src/hnsw_ace_layered_example.cu @@ -11,7 +11,10 @@ // 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. // -// Run with an optional backend (regular by default) and output directory. +// Run with an optional graph build method (regular by default) and output directory. +// regular: default hnsw::build selection; uses ordinary CAGRA when memory permits. +// ace: explicitly select ACE with disk mode disabled. +// disk-ace: explicitly select ACE with disk mode enabled in the output directory. // For example: ./hnsw_ace_layered_example regular /tmp/my_hnsw_graph // Use a fresh output directory for each disk-ace run. // @@ -92,7 +95,7 @@ auto quantize_dataset(raft::device_resources const& dev_resources, return dataset_i8; } -auto make_hnsw_params(const std::string& backend, const std::string& build_dir) +auto make_hnsw_params(const std::string& build_method, const std::string& build_dir) -> cuvs::neighbors::hnsw::index_params { using namespace cuvs::neighbors; @@ -104,10 +107,10 @@ auto make_hnsw_params(const std::string& backend, const std::string& build_dir) hnsw_params.M = 32; hnsw_params.ef_construction = 120; - if (backend != "regular") { + if (build_method != "regular") { auto ace_params = hnsw::graph_build_params::ace_params(); ace_params.npartitions = 4; - ace_params.use_disk = backend == "disk-ace"; + ace_params.use_disk = build_method == "disk-ace"; ace_params.build_dir = build_dir; hnsw_params.graph_build_params = ace_params; } @@ -124,7 +127,7 @@ auto hnsw_build(raft::device_resources const& dev_resources, using namespace cuvs::neighbors; auto hnsw_index = hnsw::build(dev_resources, hnsw_params, dataset); - // In-memory backends own a temporary artifact. Persist it before the handle is destroyed. + // In-memory build methods own a temporary artifact. Persist it before the handle is destroyed. hnsw::serialize(dev_resources, artifact_path, *hnsw_index); std::cout << " hnsw_build: layered artifact saved to " << artifact_path << std::endl; return artifact_path; @@ -192,13 +195,14 @@ void hnsw_search(raft::device_resources const& dev_resources, int main(int argc, char** argv) { - const std::string backend = argc > 1 ? argv[1] : "regular"; - const std::string build_dir = argc > 2 ? argv[2] : "/tmp/hnsw_ace_layered"; - if (argc > 3 || (backend != "regular" && backend != "ace" && backend != "disk-ace")) { + const std::string build_method = argc > 1 ? argv[1] : "regular"; + const std::string build_dir = argc > 2 ? argv[2] : "/tmp/hnsw_ace_layered"; + if (argc > 3 || + (build_method != "regular" && build_method != "ace" && build_method != "disk-ace")) { std::cerr << "Usage: " << argv[0] << " [regular|ace|disk-ace] [output-directory]" << std::endl; return 1; } - std::cout << "Graph build backend: " << backend << std::endl; + std::cout << "Graph build method: " << build_method << std::endl; raft::device_resources dev_resources; rmm::mr::pool_memory_resource pool_mr(rmm::mr::get_current_device_resource_ref(), @@ -242,7 +246,7 @@ int main(int argc, char** argv) auto dataset_i8_view = raft::make_host_matrix_view( dataset_i8.data_handle(), n_samples, n_dim); auto dataset_path = write_local_dataset(dataset_host_view, build_dir + "/dataset.npy"); - auto hnsw_params = make_hnsw_params(backend, build_dir); + auto hnsw_params = make_hnsw_params(build_method, build_dir); std::cout << "[stage 2] Build graph-only HNSW artifact from int8 data" << std::endl; auto artifact_path = hnsw_build( @@ -255,7 +259,7 @@ int main(int argc, char** argv) hnsw_search(dev_resources, *hnsw_index, queries_host_view); #else auto dataset_path = write_local_dataset(dataset_host_view, build_dir + "/dataset.npy"); - auto hnsw_params = make_hnsw_params(backend, build_dir); + auto hnsw_params = make_hnsw_params(build_method, build_dir); std::cout << "[stage 2] Build layered HNSW index" << std::endl; auto artifact_path = hnsw_build( From 0a2ad75a914fbedd94b28c7674c472ed6eafb934 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:21:27 +0200 Subject: [PATCH 12/20] Clarify serialization of loaded HNSW graphs without source files --- cpp/src/neighbors/detail/hnsw.hpp | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) diff --git a/cpp/src/neighbors/detail/hnsw.hpp b/cpp/src/neighbors/detail/hnsw.hpp index fa4e5b368b..63c8f4f85d 100644 --- a/cpp/src/neighbors/detail/hnsw.hpp +++ b/cpp/src/neighbors/detail/hnsw.hpp @@ -2630,7 +2630,18 @@ void search(raft::resources const& res, } } -// Reconstruct the graph-only artifact from a loaded index without retaining its source file. +/** + * @brief Serialize the existing graph topology of a deserialized in-memory GRAPH_ONLY index. + * + * Deserialization loads the graph and vectors into memory and closes its input files. The + * returned index retains no source graph path or file descriptor. This function writes the + * in-memory levels and links, so it still works after the original graph file has been removed. + * It does not rebuild the hierarchy or write the attached vectors. + * + * @param[in] idx A deserialized GRAPH_ONLY index with an in-memory hnswlib graph. + * @param[in] filename Destination graph-only file. An existing destination is replaced after + * the new graph has been written to a sibling temporary file. + */ template void serialize_loaded_layered_hnsw(const index_impl& idx, const std::string& filename) { From b36b15e3b3e91ff4377f697c152ec798794e1177 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:22:06 +0200 Subject: [PATCH 13/20] Name and document HNSW hierarchy construction before serialization --- cpp/src/neighbors/detail/hnsw.hpp | 27 +++++++++++++++++++++++++-- 1 file changed, 25 insertions(+), 2 deletions(-) diff --git a/cpp/src/neighbors/detail/hnsw.hpp b/cpp/src/neighbors/detail/hnsw.hpp index 63c8f4f85d..89e5feec33 100644 --- a/cpp/src/neighbors/detail/hnsw.hpp +++ b/cpp/src/neighbors/detail/hnsw.hpp @@ -1185,8 +1185,31 @@ void write_layered_base_links_from_inmem(raft::resources const& res, } } +/** + * @brief Construct an HNSW hierarchy from CAGRA and write a GRAPH_ONLY file. + * + * Reuse the CAGRA graph as the base layer, assign random HNSW levels, gather the vectors of + * promoted nodes, and build the upper-layer graphs. Write levels and adjacency lists in the + * layered format without vectors. Disk-backed ACE row IDs are mapped back to original dataset + * rows; in-memory CAGRA row IDs already refer to that original order. + * + * This is a construction operation, unlike serialize_loaded_layered_hnsw, which writes an + * existing in-memory hierarchy. Vectors are required to build upper layers even though the + * output does not contain them. + * + * @param[in] res Resources for graph construction and device transfers. + * @param[in] params HNSW construction parameters; the hierarchy must be GPU. + * @param[in] index_ In-memory CAGRA graph or disk-backed ACE graph with its row mapping. + * Composite source-index mappings are not supported. + * @param[in] dataset Optional original-order host vectors. Required for disk-backed ACE and + * indexes without an attached device dataset; otherwise the attached device vectors are used. + * @param[in] artifact_file Output graph file in an existing directory. Publication fails if + * this path already exists, preserving the existing file. + * @return Path of the completed graph-only file. Load it with a separate original-order + * dataset using the two-filename deserialize overload before searching. + */ template -auto serialize_to_layered_hnswlib( +auto build_and_serialize_layered_hnsw( raft::resources const& res, const cuvs::neighbors::hnsw::index_params& params, const CagraIndexT& index_, @@ -2396,7 +2419,7 @@ std::unique_ptr> from_cagra( directory = std::make_unique(); artifact_path = directory->path() / "hnsw_index.cuvs"; } - serialize_to_layered_hnswlib(res, params, cagra_index, dataset, artifact_path); + build_and_serialize_layered_hnsw(res, params, cagra_index, dataset, artifact_path); hnsw_index->set_file_descriptor(cuvs::util::file_descriptor(artifact_path.string(), O_RDONLY)); hnsw_index->set_artifact_directory(std::move(directory)); return hnsw_index; From bf96e23eea3aff921f3c447067b7fd25252eeb44 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:31:33 +0200 Subject: [PATCH 14/20] Clarify host vector requirements for disk-backed GRAPH_ONLY conversion --- cpp/include/cuvs/neighbors/hnsw.hpp | 14 ++++++++------ fern/pages/cpp_api/cpp-api-neighbors-hnsw.md | 4 ++-- 2 files changed, 10 insertions(+), 8 deletions(-) diff --git a/cpp/include/cuvs/neighbors/hnsw.hpp b/cpp/include/cuvs/neighbors/hnsw.hpp index fce44d6bda..dd7fb7ef1e 100644 --- a/cpp/include/cuvs/neighbors/hnsw.hpp +++ b/cpp/include/cuvs/neighbors/hnsw.hpp @@ -741,9 +741,10 @@ std::unique_ptr> from_cagra( /** * @brief Construct an hnswlib index from a host-built CAGRA index. * Requires `dataset` for in-memory indices — host builds do not store vectors in the index. - * This also applies to `GRAPH_ONLY`, which requires a GPU hierarchy. Supply vectors in original - * row order with matching row count and logical dimensions. Serialize the returned temporary - * graph artifact before destroying the handle. + * `GRAPH_ONLY` requires a GPU hierarchy and explicit `dataset` vectors for both in-memory and + * disk-backed indices. Supply vectors in original row order with matching row count and logical + * dimensions. For in-memory conversion, serialize the returned temporary graph artifact before + * destroying the handle. * * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. @@ -785,9 +786,10 @@ std::unique_ptr> from_cagra( /** * @brief Construct an hnswlib index from a host-built CAGRA index (standard dataset layout). * Requires `dataset` for in-memory indices — host builds do not store vectors in the index. - * This also applies to `GRAPH_ONLY`, which requires a GPU hierarchy. Supply vectors in original - * row order with matching row count and logical dimensions. Serialize the returned temporary - * graph artifact before destroying the handle. + * `GRAPH_ONLY` requires a GPU hierarchy and explicit `dataset` vectors for both in-memory and + * disk-backed indices. Supply vectors in original row order with matching row count and logical + * dimensions. For in-memory conversion, serialize the returned temporary graph artifact before + * destroying the handle. * * `GRAPH_ONLY` conversion does not support composite CAGRA indexes with `source_indices()` * mappings. Disk-backed ACE's internal row mapping is supported and restored to original row IDs. diff --git a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md index 471d08e5f0..962bc67709 100644 --- a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md +++ b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md @@ -522,7 +522,7 @@ Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existin **Additional overload:** `neighbors::hnsw::from_cagra` -Construct an hnswlib index from a host-built CAGRA index. Requires `dataset` for in-memory indices — host builds do not store vectors in the index. This also applies to `GRAPH_ONLY`, which requires a GPU hierarchy. Supply vectors in original row order with matching row count and logical dimensions. Serialize the returned temporary graph artifact before destroying the handle. +Construct an hnswlib index from a host-built CAGRA index. Requires `dataset` for in-memory indices — host builds do not store vectors in the index. `GRAPH_ONLY` requires a GPU hierarchy and explicit `dataset` vectors for both in-memory and disk-backed indices. Supply vectors in original row order with matching row count and logical dimensions. For in-memory conversion, serialize the returned temporary graph artifact before destroying the handle. ```cpp std::unique_ptr> from_cagra( @@ -552,7 +552,7 @@ Disk-backed `GRAPH_ONLY` conversion creates `hnsw_index.cuvs` beside the existin **Additional overload:** `neighbors::hnsw::from_cagra` -Construct an hnswlib index from a host-built CAGRA index (standard dataset layout). Requires `dataset` for in-memory indices — host builds do not store vectors in the index. This also applies to `GRAPH_ONLY`, which requires a GPU hierarchy. Supply vectors in original row order with matching row count and logical dimensions. Serialize the returned temporary graph artifact before destroying the handle. +Construct an hnswlib index from a host-built CAGRA index (standard dataset layout). Requires `dataset` for in-memory indices — host builds do not store vectors in the index. `GRAPH_ONLY` requires a GPU hierarchy and explicit `dataset` vectors for both in-memory and disk-backed indices. Supply vectors in original row order with matching row count and logical dimensions. For in-memory conversion, serialize the returned temporary graph artifact before destroying the handle. ```cpp std::unique_ptr> from_cagra( From 1b9663de96b55f24243e4f5c75d55dd071a357a4 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:31:47 +0200 Subject: [PATCH 15/20] Correct the persistent graph artifact name in the HNSW example --- examples/cpp/src/hnsw_ace_layered_example.cu | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/examples/cpp/src/hnsw_ace_layered_example.cu b/examples/cpp/src/hnsw_ace_layered_example.cu index dd55763183..65febd1e17 100644 --- a/examples/cpp/src/hnsw_ace_layered_example.cu +++ b/examples/cpp/src/hnsw_ace_layered_example.cu @@ -8,8 +8,9 @@ // 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. +// The build creates hnsw_index.cuvs; this example serializes a persistent copy to saved_graph.cuvs. +// Transfer saved_graph.cuvs to the search server. The dataset remains separate and does not +// need to be transferred when the search server already has it locally. // // Run with an optional graph build method (regular by default) and output directory. // regular: default hnsw::build selection; uses ordinary CAGRA when memory permits. @@ -27,7 +28,7 @@ // // Layered-on-disk layout: // -// index_dir/hnsw_index.cuvs +// index_dir/saved_graph.cuvs // fixed header + layer descriptors // levels: uint8 [N], max HNSW level for each original row id // base nodes + base links: uint32 node ids with hnswlib-ready link rows From 361b8caea0de6c8d1d4d359fded40ff98b218c86 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:32:16 +0200 Subject: [PATCH 16/20] Align the layered HNSW NVTX range with hierarchy construction --- cpp/src/neighbors/detail/hnsw.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/src/neighbors/detail/hnsw.hpp b/cpp/src/neighbors/detail/hnsw.hpp index 89e5feec33..a047b02b9d 100644 --- a/cpp/src/neighbors/detail/hnsw.hpp +++ b/cpp/src/neighbors/detail/hnsw.hpp @@ -1219,7 +1219,8 @@ auto build_and_serialize_layered_hnsw( using IdxT = typename CagraIndexT::index_type; static_assert(is_cagra_hnsw_export_index_v, "Layered HNSW serialization expects a CAGRA index type"); - raft::common::nvtx::range fun_scope("hnsw::serialize_layered"); + raft::common::nvtx::range fun_scope( + "hnsw::build_and_serialize_layered_hnsw"); const auto total_start_time = std::chrono::steady_clock::now(); const bool disk_graph = index_.graph_fd().has_value(); From 7eebd1546bd9f6c0debf409fa20eb0e7ed558504 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Thu, 15 Oct 2026 16:32:35 +0200 Subject: [PATCH 17/20] Describe artifact-backed HNSW serialization independently of its build source --- cpp/src/neighbors/detail/hnsw.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/neighbors/detail/hnsw.hpp b/cpp/src/neighbors/detail/hnsw.hpp index a047b02b9d..96ade71524 100644 --- a/cpp/src/neighbors/detail/hnsw.hpp +++ b/cpp/src/neighbors/detail/hnsw.hpp @@ -2785,9 +2785,9 @@ void serialize(raft::resources const& res, const std::string& filename, const in return; } - // Check if this is a disk-based index (created from disk-backed CAGRA) + // Check whether the index is backed by an existing artifact file. if (idx_impl && idx_impl->file_descriptor().has_value()) { - // For disk-based indexes, copy the existing file to the new location + // Copy the existing artifact, including temporary output from in-memory builds. std::string source_path = idx_impl->file_path(); RAFT_EXPECTS(!source_path.empty(), "Disk-based index has invalid file path"); RAFT_EXPECTS(std::filesystem::exists(source_path), From c2e8fe4b9ed8c22fbea725aeb2bff71eb6f437f0 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Fri, 16 Oct 2026 13:27:39 +0200 Subject: [PATCH 18/20] Reuse RAFT transfer helpers in the layered HNSW gather --- cpp/src/neighbors/hnsw_layered_gather.cu | 15 +++++---------- 1 file changed, 5 insertions(+), 10 deletions(-) diff --git a/cpp/src/neighbors/hnsw_layered_gather.cu b/cpp/src/neighbors/hnsw_layered_gather.cu index bfa4822e86..0c17c042a0 100644 --- a/cpp/src/neighbors/hnsw_layered_gather.cu +++ b/cpp/src/neighbors/hnsw_layered_gather.cu @@ -50,19 +50,14 @@ void gather_layered_hnsw_vectors(raft::resources const& res, rmm::device_uvector values(batch_rows * row_bytes, stream); for (size_t start = 0; start < rows; start += batch_rows) { const auto count = std::min(batch_rows, rows - start); - RAFT_CUDA_TRY(cudaMemcpyAsync(row_ids.data(), - host_row_ids + start, - count * sizeof(size_t), - cudaMemcpyHostToDevice, - stream)); + raft::update_device(row_ids.data(), host_row_ids + start, count, stream); raft::matrix::gather( source_bytes, row_bytes, source_rows, row_ids.data(), count, values.data(), stream); RAFT_CUDA_TRY(cudaGetLastError()); - RAFT_CUDA_TRY(cudaMemcpyAsync(static_cast(host_output) + start * row_bytes, - values.data(), - count * row_bytes, - cudaMemcpyDeviceToHost, - stream)); + raft::update_host(static_cast(host_output) + start * row_bytes, + values.data(), + count * row_bytes, + stream); raft::resource::sync_stream(res); } } From ad46b4cba7434376bd0cc32f94f40dea2571da25 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Fri, 16 Oct 2026 13:27:54 +0200 Subject: [PATCH 19/20] Share layered HNSW descriptor encoding between graph writers --- cpp/src/neighbors/detail/hnsw.hpp | 34 ++++++++++++++++--------------- 1 file changed, 18 insertions(+), 16 deletions(-) diff --git a/cpp/src/neighbors/detail/hnsw.hpp b/cpp/src/neighbors/detail/hnsw.hpp index 96ade71524..2fdecc572d 100644 --- a/cpp/src/neighbors/detail/hnsw.hpp +++ b/cpp/src/neighbors/detail/hnsw.hpp @@ -792,6 +792,21 @@ inline auto metric_name(cuvs::distance::DistanceType metric) -> const char* } } +inline auto make_layered_hnsw_descriptors(const layered_hnsw_file_metadata& metadata) + -> std::vector +{ + std::vector layer_descriptors(metadata.layers.size()); + for (size_t i = 0; i < metadata.layers.size(); ++i) { + const auto& layer = metadata.layers[i]; + layer_descriptors[i] = {static_cast(layer.level), + static_cast(layer.row_count), + static_cast(layer.degree), + static_cast(layer.node_offset), + static_cast(layer.link_offset)}; + } + return layer_descriptors; +} + template auto make_layered_hnsw_header(const layered_hnsw_file_metadata& metadata, cuvs::distance::DistanceType metric) -> layered_hnsw_file_header @@ -1300,16 +1315,8 @@ auto build_and_serialize_layered_hnsw( metadata.upper_nodes_bytes = metadata.upper_nodes_count * sizeof(IdxT); metadata.upper_links_bytes = next_upper_link_offset * metadata.upper_link_row_bytes; - const auto header = make_layered_hnsw_header(metadata, index_.metric()); - std::vector layer_descriptors(metadata.layers.size()); - for (size_t i = 0; i < metadata.layers.size(); ++i) { - const auto& layer = metadata.layers[i]; - layer_descriptors[i] = {static_cast(layer.level), - static_cast(layer.row_count), - static_cast(layer.degree), - static_cast(layer.node_offset), - static_cast(layer.link_offset)}; - } + const auto header = make_layered_hnsw_header(metadata, index_.metric()); + const auto layer_descriptors = make_layered_hnsw_descriptors(metadata); const auto descriptors_offset = sizeof(layered_hnsw_file_header); const auto descriptors_bytes = layer_descriptors.size() * sizeof(layered_hnsw_layer_descriptor); @@ -2695,7 +2702,6 @@ void serialize_loaded_layered_hnsw(const index_impl& idx, const std::string& metadata.base_degree = std::max( metadata.base_degree, static_cast(graph->getListCount(graph->get_linklist0(row)))); } - std::vector descriptors; for (int level = 1; level <= metadata.maxlevel; ++level) { const auto count = static_cast(std::count_if( levels.begin(), levels.end(), [level](uint8_t value) { return value >= level; })); @@ -2705,17 +2711,13 @@ void serialize_loaded_layered_hnsw(const index_impl& idx, const std::string& degree, metadata.upper_nodes_count, metadata.upper_nodes_count}); - descriptors.push_back({static_cast(level), - count, - degree, - metadata.upper_nodes_count, - metadata.upper_nodes_count}); metadata.upper_nodes_count += count; } metadata.upper_nodes_bytes = metadata.upper_nodes_count * sizeof(uint32_t); metadata.upper_links_bytes = metadata.upper_nodes_count * metadata.upper_link_row_bytes; auto header = make_layered_hnsw_header(metadata, idx.metric()); header.construction_dtype = static_cast(idx.construction_dtype()); + const auto descriptors = make_layered_hnsw_descriptors(metadata); const size_t descriptors_bytes = descriptors.size() * sizeof(layered_hnsw_layer_descriptor); const size_t levels_offset = align_up(sizeof(header) + descriptors_bytes, layered_hnsw_alignment); const size_t base_nodes_offset = levels_offset + metadata.levels_bytes; From 3bbca2fb6bc8511ce4559c002e52064ad7bc8f39 Mon Sep 17 00:00:00 2001 From: Julian Miller Date: Fri, 16 Oct 2026 13:28:05 +0200 Subject: [PATCH 20/20] Share the float and quantized HNSW example load and search flow --- examples/cpp/src/hnsw_ace_layered_example.cu | 19 +++++-------------- 1 file changed, 5 insertions(+), 14 deletions(-) diff --git a/examples/cpp/src/hnsw_ace_layered_example.cu b/examples/cpp/src/hnsw_ace_layered_example.cu index 65febd1e17..a898822223 100644 --- a/examples/cpp/src/hnsw_ace_layered_example.cu +++ b/examples/cpp/src/hnsw_ace_layered_example.cu @@ -241,37 +241,28 @@ int main(int argc, char** argv) queries_host.data_handle(), n_queries, n_dim); std::filesystem::create_directories(build_dir); + auto dataset_path = write_local_dataset(dataset_host_view, build_dir + "/dataset.npy"); + auto hnsw_params = make_hnsw_params(build_method, build_dir); #if HNSW_ACE_LAYERED_USE_QUANTIZATION auto dataset_i8 = quantize_dataset(dev_resources, dataset_host_view); auto dataset_i8_view = raft::make_host_matrix_view( dataset_i8.data_handle(), n_samples, n_dim); - auto dataset_path = write_local_dataset(dataset_host_view, build_dir + "/dataset.npy"); - auto hnsw_params = make_hnsw_params(build_method, build_dir); std::cout << "[stage 2] Build graph-only HNSW artifact from int8 data" << std::endl; auto artifact_path = hnsw_build( dev_resources, hnsw_params, dataset_i8_view, build_dir + "/saved_graph.cuvs"); - - std::cout << "[stage 3] Attach original float dataset" << std::endl; - auto hnsw_index = hnsw_deserialize(dev_resources, artifact_path, dataset_path); - - std::cout << "[stage 4] Search float HNSW index" << std::endl; - hnsw_search(dev_resources, *hnsw_index, queries_host_view); #else - auto dataset_path = write_local_dataset(dataset_host_view, build_dir + "/dataset.npy"); - auto hnsw_params = make_hnsw_params(build_method, build_dir); - std::cout << "[stage 2] Build layered HNSW index" << std::endl; auto artifact_path = hnsw_build( dev_resources, hnsw_params, dataset_host_view, build_dir + "/saved_graph.cuvs"); +#endif - std::cout << "[stage 3] Deserialize layered HNSW index" << std::endl; + std::cout << "[stage 3] Attach original float dataset" << std::endl; auto hnsw_index = hnsw_deserialize(dev_resources, artifact_path, dataset_path); - std::cout << "[stage 4] Search HNSW index" << std::endl; + std::cout << "[stage 4] Search float HNSW index" << std::endl; hnsw_search(dev_resources, *hnsw_index, queries_host_view); -#endif return 0; }