diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 9f9763160c..cb165fe982 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -1484,6 +1484,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..dd7fb7ef1e 100644 --- a/cpp/include/cuvs/neighbors/hnsw.hpp +++ b/cpp/include/cuvs/neighbors/hnsw.hpp @@ -45,9 +45,18 @@ 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. + * + * 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, @@ -64,8 +73,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 +207,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 ""; } @@ -486,12 +498,24 @@ 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. + * + * 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 * @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 +550,24 @@ 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. + * + * 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 * @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 +602,24 @@ 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. + * + * 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 * @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 +654,24 @@ 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. + * + * 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 * @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 +698,17 @@ 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. + * + * 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, @@ -671,6 +741,19 @@ 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. + * `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. + * + * 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, @@ -703,6 +786,19 @@ 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. + * `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. + * + * 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, @@ -1103,7 +1199,12 @@ 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. + * 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 @@ -1134,7 +1235,12 @@ 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. + * 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 @@ -1165,7 +1271,12 @@ 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. + * 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 @@ -1196,7 +1307,12 @@ 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. + * 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/cpp/src/neighbors/detail/hnsw.hpp b/cpp/src/neighbors/detail/hnsw.hpp index 19c0afd820..2fdecc572d 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 @@ -750,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 @@ -1073,42 +1130,147 @@ 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); + } +} + +/** + * @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 build_and_serialize_layered_hnsw( 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, "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(); - 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"); @@ -1153,16 +1315,8 @@ auto serialize_to_layered_hnswlib_from_disk( 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); @@ -1199,12 +1353,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 +1383,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 +2405,34 @@ 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"); + // 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(); + artifact_path = directory->path() / "hnsw_index.cuvs"; + } + 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; + } + // 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 +2448,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 +2467,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,14 +2661,135 @@ void search(raft::resources const& res, } } +/** + * @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) +{ + 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)))); + } + 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}); + 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; + 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); - // Check if this is a disk-based index (created from disk-backed CAGRA) + if (idx_impl && idx_impl->output_format() == HnswOutputFormat::GRAPH_ONLY && + idx_impl->get_index() != nullptr) { + serialize_loaded_layered_hnsw(*idx_impl, filename); + return; + } + + // 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), @@ -2878,6 +3182,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 +3245,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..a07f1bab92 --- /dev/null +++ b/cpp/src/neighbors/detail/hnsw_layered_gather.hpp @@ -0,0 +1,41 @@ +/* + * 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 { +/** + * @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, + 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..0c17c042a0 --- /dev/null +++ b/cpp/src/neighbors/hnsw_layered_gather.cu @@ -0,0 +1,64 @@ +/* + * 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 +#include +#include +#include + +namespace cuvs::neighbors::hnsw::detail { +namespace { +// 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, + 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"); + 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); + 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::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::update_host(static_cast(host_output) + start * row_bytes, + values.data(), + count * row_bytes, + stream); + raft::resource::sync_stream(res); + } +} +} // namespace cuvs::neighbors::hnsw::detail 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 c50d0c8cee..77ffd9b90b 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -355,7 +355,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 ) @@ -486,7 +486,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..c0921b92f8 100644 --- a/cpp/tests/neighbors/ann_hnsw_ace.cuh +++ b/cpp/tests/neighbors/ann_hnsw_ace.cuh @@ -33,6 +33,23 @@ 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, // 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 { int n_queries; int n_rows; @@ -444,7 +461,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 +520,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 +534,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 +597,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 +671,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 +716,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..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(); @@ -74,4 +117,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/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 +96,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& build_method, const std::string& build_dir) + -> cuvs::neighbors::hnsw::index_params { using namespace cuvs::neighbors; @@ -101,11 +108,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 (build_method != "regular") { + auto ace_params = hnsw::graph_build_params::ace_params(); + ace_params.npartitions = 4; + ace_params.use_disk = build_method == "disk-ace"; + ace_params.build_dir = build_dir; + hnsw_params.graph_build_params = ace_params; + } return hnsw_params; } @@ -113,16 +122,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 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; } @@ -186,8 +194,16 @@ void hnsw_search(raft::device_resources const& dev_resources, } // namespace -int main() +int main(int argc, char** argv) { + 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 method: " << build_method << std::endl; raft::device_resources dev_resources; rmm::mr::pool_memory_resource pool_mr(rmm::mr::get_current_device_resource_ref(), @@ -224,38 +240,29 @@ 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); + 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, std::string{kBuildDir} + "/dataset.npy"); - auto hnsw_params = make_hnsw_ace_params(kBuildDir); - std::cout << "[stage 2] Build graph-only HNSW artifact from int8 data with ACE" << std::endl; - auto artifact_path = hnsw_build(dev_resources, hnsw_params, dataset_i8_view); + std::cout << "[stage 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"); +#else + 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] Attach original float dataset" << std::endl; auto hnsw_index = hnsw_deserialize(dev_resources, artifact_path, dataset_path); std::cout << "[stage 4] Search float HNSW index" << std::endl; hnsw_search(dev_resources, *hnsw_index, queries_host_view); -#else - auto dataset_path = - write_local_dataset(dataset_host_view, std::string{kBuildDir} + "/dataset.npy"); - auto hnsw_params = make_hnsw_ace_params(kBuildDir); - - std::cout << "[stage 2] Build layered HNSW index with ACE" << std::endl; - auto artifact_path = hnsw_build(dev_resources, hnsw_params, dataset_host_view); - - std::cout << "[stage 3] Deserialize layered HNSW index" << std::endl; - auto hnsw_index = hnsw_deserialize(dev_resources, artifact_path, dataset_path); - - std::cout << "[stage 4] Search HNSW index" << std::endl; - hnsw_search(dev_resources, *hnsw_index, queries_host_view); -#endif return 0; } diff --git a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md index b5114bc4ac..43a91b0dbd 100644 --- a/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md +++ b/fern/pages/cpp_api/cpp-api-neighbors-hnsw.md @@ -36,7 +36,11 @@ 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. + +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 { @@ -71,7 +75,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 +189,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` @@ -344,7 +350,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/.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 +361,10 @@ 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. + +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: **Parameters** @@ -364,7 +374,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 +386,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/.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 +397,10 @@ 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. + +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: **Parameters** @@ -396,7 +410,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 +422,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/.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 +433,10 @@ 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. + +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: **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/.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,10 @@ 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. + +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: **Parameters** @@ -460,7 +482,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 +501,11 @@ 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. + +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** @@ -496,7 +522,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. `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( @@ -507,6 +533,10 @@ 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. + +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** | Name | Direction | Type | Description | @@ -522,7 +552,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. `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( @@ -533,6 +563,10 @@ 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. + +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** | Name | Direction | Type | Description | @@ -802,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. +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); @@ -824,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. +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); @@ -846,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. +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); @@ -868,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. +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);