Add General HNSW GRAPH_ONLY Support

cpp/CMakeLists.txt

@@ -1484,6 +1484,7 @@ if(NOT BUILD_CPU_ONLY)
     src/neighbors/composite/index.cu
     $<$<BOOL:${BUILD_CAGRA_HNSWLIB}>:src/neighbors/cagra.cpp>
     $<$<BOOL:${BUILD_CAGRA_HNSWLIB}>:src/neighbors/hnsw.cpp>
+    $<$<BOOL:${BUILD_CAGRA_HNSWLIB}>:src/neighbors/hnsw_layered_gather.cu>
     src/neighbors/ivf_common.cu
     src/neighbors/ivf_flat_index.cpp
     ${ivf_flat_build_extend_inst_files}

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 <cstddef>
+#include <cstdint>
+#include <raft/core/resources.hpp>
+
+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

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 <algorithm>
+#include <raft/core/error.hpp>
+#include <raft/core/resource/cuda_stream.hpp>
+#include <raft/matrix/gather.cuh>
+#include <raft/util/cudart_utils.hpp>
+#include <rmm/device_uvector.hpp>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/transform_iterator.h>
+
+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<size_t>(0),
+    strided_source_byte{static_cast<const uint8_t*>(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<size_t>(1, (64 * 1024 * 1024) / (row_bytes + sizeof(size_t))));
+  rmm::device_uvector<size_t> row_ids(batch_rows, stream);
+  rmm::device_uvector<uint8_t> 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<uint8_t*>(host_output) + start * row_bytes,
+                      values.data(),
+                      count * row_bytes,
+                      stream);
+    raft::resource::sync_stream(res);
+  }
+}
+}  // namespace cuvs::neighbors::hnsw::detail

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<T, half>) {
+    // Older cuVS writers used RAFT's legacy half descriptor instead of NumPy's "<f2".
+    constexpr char legacy_half_dtype[] = "<e2";
+    if (std::equal(dtype_prefix, dtype_prefix + dtype_prefix_size, legacy_half_dtype)) {
+      return true;
+    }
+  }
+
   auto expected_dtype = cuvs::util::detail::numpy_dtype_string<T>();
   expected_dtype.resize(dtype_prefix_size, '\0');
 

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
 )

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<AnnHnswAceInputs> {
     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<IdxT> indexes_naive(queries_size);
@@ -503,22 +520,71 @@ class AnnHnswAceTest : public ::testing::TestWithParam<AnnHnswAceInputs> {
     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;
     EXPECT_THROW(
       hnsw::build(handle_, invalid_params, raft::make_const_mdspan(database_host.view())),
       std::exception);
 
-    auto hnsw_index =
-      hnsw::build(handle_, hnsw_params, raft::make_const_mdspan(database_host.view()));
+    std::unique_ptr<hnsw::index<DataT>> 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<const DataT, int64_t>(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<DataT> 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<const DataT, int64_t>(
+          database_host.data_handle(), ps.n_rows - 1, ps.dim);
+        auto wrong_columns = raft::make_host_matrix_view<const DataT, int64_t>(
+          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<AnnHnswAceInputs> {
     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<uint64_t, int64_t>(ps.n_queries, ps.k);
@@ -605,6 +671,17 @@ class AnnHnswAceTest : public ::testing::TestWithParam<AnnHnswAceInputs> {
       hnsw::deserialize(handle_, artifact_path, float_dataset_file, &mixed_precision_index);
       ASSERT_NE(mixed_precision_index, nullptr);
       std::unique_ptr<hnsw::index<float>> 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<char*>(&mixed_header), sizeof(mixed_header));
+      ASSERT_TRUE(mixed_file.good());
+      EXPECT_EQ(mixed_header.construction_dtype,
+                static_cast<uint32_t>(cuvs::neighbors::hnsw::detail::layered_hnsw_dtype::int8));
+      hnsw::index<float>* 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<AnnHnswAceInputs> {
     }
     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<DataT>* copied_index = nullptr;
     hnsw::deserialize(handle_, copied_artifact, dataset_file, &copied_index);
     ASSERT_NE(copied_index, nullptr);
     std::unique_ptr<hnsw::index<DataT>> copied_guard(copied_index);
+    hnsw::serialize(handle_, copied_artifact, *copied_guard);
+    hnsw::index<DataT>* 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,