feat: GPU_HNSW / GPU_HNSW_SQ index using faiss::gpu::GpuIndexHNSW — native FP32/FP16/BF16/INT8

[devin-ai-integration]

Summary

GPU HNSW index implementation using a vendored FAISS GpuIndexHNSW backend. Accepts CPU-built HNSW/HNSW_SQ indexes at load time, uploads vectors and graph to GPU VRAM, frees the CPU copy, and runs a parallel beam search entirely on GPU with per-segment CUDA streams. Recall is tuned via ef.

Registered index types & precisions: GPU_HNSW and GPU_HNSW_SQ, each for fp32, fp16, bf16, and int8. GPU_HNSW_SQ is a thin alias — GpuHnswSQIndexNode subclasses GpuHnswIndexNode and only overrides Type(); the GPU search and load paths are identical.

Registrations (both index types) now cover all four precisions:

// index_table.h — legal (index, VecType) pairs
{GPU_HNSW, VECTOR_FLOAT}, {GPU_HNSW, VECTOR_FLOAT16},
{GPU_HNSW, VECTOR_BFLOAT16}, {GPU_HNSW, VECTOR_INT8},   // + same 4 for GPU_HNSW_SQ

// faiss_hnsw.cc
KNOWHERE_REGISTER_GLOBAL(GPU_HNSW, ...<fp16>..., fp16, true, feature::FP16 | feature::GPU);
KNOWHERE_REGISTER_GLOBAL(GPU_HNSW, ...<bf16>..., bf16, true, feature::BF16 | feature::GPU);
IndexStaticFaced<fp16>::Instance().RegisterStaticFunc<GpuHnswIndexNode>(INDEX_GPU_HNSW);
IndexStaticFaced<bf16>::Instance().RegisterStaticFunc<GpuHnswIndexNode>(INDEX_GPU_HNSW);
// (+ mirror for GPU_HNSW_SQ)

Native low-precision storage: fp16/bf16/int8 stay in their on-disk layout on the device (2/2/1 bytes/element) and are up-converted to fp32 per element inside the search kernel — mirroring the pre-existing INT8 path. The vendored FAISS copy (thirdparty/faiss/faiss/gpu/impl/GpuHnsw*.{h,cuh}) is kept byte-identical to 6si/faiss#1: a GpuHnswDatasetType enum drives kernel dispatch and upload_halfwidth_dataset() copies the raw 2-byte codes verbatim. The IndexNodeDataMockWrapper<fp16|bf16> converts fp16/bf16 query inputs to the node's fp32-facing interface, while the stored dataset stays native low-precision.

Load-path fix (mmap/file) — fixes the production querynode OOM: the eager GPU upload must run on both deserialization entrypoints. Milvus loads sealed segments through the mmap/file path (VectorMemIndex::LoadFromFile → index_.DeserializeFromFile()), not Deserialize(BinarySet). GpuHnswIndexNode previously only overrode Deserialize(BinarySet), so DeserializeFromFile fell through to the base CPU loader and the upload never ran: the CPU-built HNSW stayed resident in host RAM (fp32-expanded for int8/fp16/bf16 inputs, ~5.9 GiB/segment for mpd_v2) while VRAM sat idle, OOMKilling querynodes during the load storm before any search could trigger the lazy-upload fallback. Fix factors the upload/unpublish into helpers and calls them from both entrypoints:

Status DeserializeFromFile(filename, config) override {
    std::unique_lock lock(gpu_mutex_);
    UnpublishGpuIndexLocked();
    auto st = BaseFaissRegularIndexHNSWNode::DeserializeFromFile(filename, config);  // load CPU index
    if (st != Status::success) return st;
    return UploadCpuIndexToGpuLocked();  // copyFromWithMetric -> free indexes[0] -> publish gpu_index_
}

GpuHnswSQIndexNode inherits the fix. New [gpu_hnsw_p1] test loads via DeserializeFromFile (enable_mmap true/false) and asserts Size()==0 before any search (i.e. indexes[0] freed) — fails pre-fix and cannot be masked by the first-search lazy upload.

Architecture (unchanged core):

• GpuHnswIndexNode wraps faiss::gpu::GpuIndexHNSW; accepts CPU-serialized HNSW/HNSW_SQ at Deserialize() (binset key aliased from GPU_HNSW_SQ/HNSW_SQ/HNSW) and at DeserializeFromFile() (raw faiss index bytes / mmap). CPU copy freed after upload; Count()/Dim() captured before the free.
• Shared StandardGpuResources singleton (GetSharedGpuResources()); StaticEstimateLoadResource returns memoryCost=0; GetGpuConstructionMutex() serializes construction; gpu_ready_ publishes the built index for the lock-free search fast path; per-segment cudaStreamNonBlocking streams.

Tests: added Test GPU HNSW FP16 Deserialization and Test GPU HNSW BF16 Deserialization (CPU HNSW_SQ QT_fp16/QT_bf16 build → GPU load → recall vs brute force; floors 0.9 / 0.85), plus the [gpu_hnsw_p1] regression sections (cosine-norm parity, reload lifetime, load accounting, and the new DeserializeFromFile upload/free assertion). Guarded by KNOWHERE_WITH_CUVS, so they compile/run only on GPU builds (no silent no-op pass on CPU builds).

Kernel: 1-thread-per-distance, 128 concurrent candidates/block, num_parents == 0 stagnation break. The earlier non-functional "OCQ" overflow queue was removed here and in faiss#1 (decision record docs/design-docs/design_docs/20260711-gpu-hnsw-ocq-removal.md in 6si/milvus#6); use a larger ef for higher recall.

Filtered search: Search() rejects a non-empty bitset with invalid_args (delete/TTL/partition). No in-node CPU fallback (copy freed after upload); GPU-side bitset filtering is a documented follow-up. Scoped to append-mostly / immutable collections for now.

Verification: compile-verified via the GPU build. The DeserializeFromFile fix is validated by the [gpu_hnsw_p1] unit test and a live single-segment + full mpd_v2 load re-run (VRAM growth + host-anon-returns-to-baseline + zero OOMKills) — in progress. The prior perf numbers (8× g7e.2xlarge, mpd_v2 569M, GPU 3.3–4.5× CPU vec/s, R@1 0.892) predate this change and covered fp32/int8; native fp16/bf16 correctness is validated by the GPU unit/recall gates and the sealed-segment cluster e2e (see Performance below); a matched fp16/bf16 throughput run is a follow-up.

Performance (CPU-HNSW vs GPU)

Benchmarked on the same dataset and harness (collection mpd_v2, 569.9M rows, INT8 COSINE; run_eval_jobs.py, --sample-size 2048 --continuous-iterations 30). Best cell = batch 512, 8 parallel × 4 workers (32 concurrency):

Cluster	Nodes	Throughput	p50 latency	R@1
GPU HNSW (g7e.2xlarge)	8	28,087 vec/s	518 ms	0.897
CPU HNSW (r8g.2xlarge Graviton)	32	4,365 vec/s	3,683 ms	0.889
CPU HNSW (r8g.2xlarge Graviton)	10	1,509 vec/s	9,178 ms	0.891

• GPU on 8 nodes ≈ 6.4× the throughput of CPU on 32 nodes and ≈ 18.6× the 10-node CPU cluster, at ~7× lower p50 latency.
• Peak GPU throughput = 30,913 vec/s (batch 256, 8 parallel × 1 worker; p50 62 ms, R@1 0.897).
• Recall parity throughout: R@1 ≈ 0.889–0.900 on both CPU and GPU.
• CPU HNSW is CPU-bound — each node pins at its 8-vCPU limit; throughput scales ~linearly with node count while extra concurrency only inflates latency (CPU 8×4 b1024 → p50 18s, p99 30s). GPU sweet spot is batch 256–512; batch 1024 only helps at the highest concurrency.

Throughput figures are from the gpu-hnsw-v89 image (this beam-search kernel; fp32/int8 COSINE). v89 is 10–24% faster than v83 at all matched batch-512 cells (4×1 +11%, 4×4 +11%, 8×1 +24%, 8×4 +14%), so the filtered-search path does not regress unfiltered throughput. Native fp16/bf16 throughput is a follow-up; correctness across fp32/fp16/bf16/int8 × L2/IP/COSINE (and filtered deletes/TTL/partition at CPU parity) is covered by the unit/recall/sanitizer gates and the sealed-segment cluster e2e.

Related PRs: 6si/faiss#1 (FAISS kernel), 6si/milvus#6 (Milvus integration).

Link to Devin session: https://6sense.devinenterprise.com/sessions/d39aba56b8a3467cbbf231ab631a06ed

Link to Devin session: https://6sense.devinenterprise.com/sessions/55dab0bd33c346df99b5818b30059342
Requested by: @premal

[devin-ai-integration]

🤖 Devin AI Engineer

I'll be helping with this pull request! Here's what you should know:

✅ I will automatically:

• Address comments on this PR. Add '(aside)' to your comment to have me ignore it.
• Look at CI failures and help fix them

Note: I can only respond to comments from users who have write access to this repository.

⚙️ Control Options:

• Disable automatic comment, CI, and merge conflict monitoring

[ox-security]

Successfully scanned changes introduced in a pull request into main from gpu-hnsw-faiss.

Internal scan identifier: 04b82793-19d9-4c45-a031-92521fda40d1.

Total issues	Blocking issues	Scan status
1	0	✔️

Category	Issues
Code Security	1

See all issues found during this scan in the OX Security Application.

Detailed information

Issue #	1
Name	Unsafe Func • C - Testing Code • Taint Flow • Public Repo
Status	New
Enforcement	Monitor
Severity	Low
Category	Code Security
Source tools	OX Code Security
Recommendation	Validate the size of parameters at the function entry and use bounded counterparts (strncpy, strlcpy, snprintf, memcpy_s) instead of the unsafe APIs.

3 aggregations

File	Match
thirdparty/faiss/tests/test_read_index_deserialize.cpp	memcpy(writer.data.data() + offset, &corrupt_cs, sizeof(size_t));
thirdparty/faiss/tests/test_read_index_deserialize.cpp	memcpy(&buf[i + sizeof(uint8_t)], &bbs, sizeof(int));
thirdparty/faiss/tests/test_read_index_deserialize.cpp	memcpy(&buf[i + sizeof(int)], &qbs, sizeof(int));

[devin-ai-integration]

Devin Review found 1 new potential issue.

View 5 additional findings in Devin Review.

[premal]

Code Review — Delta (INT8 DP4A + Padded Staging)

⚠️ Critical — Still Not Fixed

[C1] Lazy-upload path in Search() uses config-based metric detection — can produce wrong results for cosine/IP

The eager upload path (UploadCpuIndexToGpuLocked) correctly detects cosine/IP via dynamic_cast on the faiss index:

bool is_cosine =
    dynamic_cast<const HasInverseL2Norms*>(faiss_idx) != nullptr;
bool use_ip = is_cosine || (faiss_idx->metric_type == METRIC_INNER_PRODUCT);

But the lazy-upload fallback in Search() still reads from config:

bool is_cosine = IsMetricType(hnsw_cfg.metric_type.value(), metric::COSINE);
bool use_ip = IsMetricType(hnsw_cfg.metric_type.value(), metric::IP) || is_cosine;

If metric_type is absent from the config (empty json defaults to L2), a cosine/IP index will silently search with L2 distance → wrong results. This is the same class of bug the eager path was specifically hardened against.

Fix: Factor metric detection into a shared helper, or have the lazy path call UploadCpuIndexToGpuLocked() directly (it already handles null-index no-op and is guarded by gpu_mutex_, which the lazy path holds).

Important

[D1] INT8 registration comments still say "+128 shift" — Both int8 registration lambdas say searchHostInt8() "applies the +128 shift" but it does not (queries uploaded verbatim as signed int8). The GpuIndexHNSW.h header docstring was corrected in this delta, but these registration comments were not updated.

[M6] Only ef is threaded from knowhere config — search_width, max_iterations, and thread_block_size always use defaults. Users cannot tune these through knowhere. Acceptable for initial integration, but worth documenting.

New Changes — Verified Correct

• DP4A cosine post-processing is correct — divides distances by query norm, negates to positive. Database norms applied in-kernel. No ids/dists swap bug.
• dim % 4 == 0 guard for DP4A is a bug fix — old code had no guard and would corrupt on non-multiple-of-4 dimensions. Fallback to fp32 path is correct.
• Parallel bitonic merge (fixes previous I1 — single-threaded merge bottleneck) — implementation is correct: thread ownership is exclusive (thread i writes only slot i), negative-index guards prevent speculative OOB reads, tie-breaking is consistent.
• Unified <DataT, QueryT, USE_DP4A> template eliminates ~90 lines of duplication (fixes previous I3).
• Power-of-2 padded staging for non-power-of-two 2*M graphs — correctly handled with sentinels.

Previous Findings Status

Finding	Status
C1 (lazy-upload metric detection)	❌ Not fixed
I1 (single-threaded merge)	✅ Fixed — parallel bitonic merge
I2 (stale "+128 shift" docstrings)	⚠️ Partially fixed — header fixed, registration comments stale
I3 (duplicated upper-layer logic)	✅ Fixed — unified template
I4 (maxMemoryCost under-counting)	📄 Documented — detailed limitation comment added

Full review: static analysis only (no CUDA/GPU available for compile/test verification).

[premal]

Code Review — Delta 2 (Bitonic Sort Race Fix + Recall Tests)

Bug fixes correctly mirror faiss PR

The vendored kernel changes are identical to 6si/faiss#1: two-phase barrier bitonic sort fix, ri_mid = pos - mid - 1 off-by-one fix, and defensive guards (entry-point, parent-ID, result-ID validation). All verified correct via brute-force simulation (see faiss PR comment).

⚠️ C1 (Critical) — STILL NOT FIXED (3rd review, unchanged)

The lazy-upload path in Search() still uses config-based metric_type detection instead of dynamic_cast. If metric_type is absent from config (defaults to L2), a cosine/IP index silently searches with L2 → wrong results. This has been flagged in every review cycle and remains the single most important outstanding fix.

New recall tests — well-designed

• 5 dtypes x 3 metrics = 15 cases, one TEST_CASE per dtype for process isolation
• CPU-vs-GPU recall against brute-force ground truth (correct reference)
• Thresholds: fp32/fp16 >= 0.85, bf16 >= 0.80, sq8 >= 0.65, int8 >= 0.85 — each justified by dtype precision characteristics
• Not too loose (catches regressions), not too tight (won't flake on quantization variance)
• Good design: fp32 as control, dim=64 for DP4A coverage, fixed seed

Still Open (carried forward)

• C1: Lazy-upload metric detection — not fixed
• D1: INT8 registration comments still say "+128 shift"
• M6: Only ef threaded from config
• I4: maxMemoryCost under-counting for fp32/SQ8 (documented, not fixed)

No new critical issues in this delta. The fixes address real bugs (50% merge mis-partition rate + shared-memory race).

[premal]

Code Review — Delta 3 (C1 fix, distance convention, SQ alias, recall parity)

✅ C1 (Critical) — FINALLY FIXED

The lazy-upload path in Search() now calls UploadCpuIndexToGpuLocked() directly, making metric detection identical to the eager paths (Deserialize/DeserializeFromFile). All three upload paths now share the same dynamic_cast-based metric detection. The const_cast on indexes[0] is documented. This was the single most important outstanding fix — flagged in every previous review cycle.

New Issue — Important (lower severity than C1)

[I1] Post-processing still reads hnsw_cfg.metric_type.value() for cosine detection — The query-side normalization (fp32 path: normalize queries; int8 path: scale distances by 1/||q||) still reads metric_type from the search config. If metric_type is missing from config, cosine queries won't be normalized → wrong results. This is a residual of the same C1 bug class but lower severity: it only affects query-side normalization (not graph traversal), and only triggers when config omits metric_type (Milvus always provides it). Consider deriving the metric from gpu_index_ (which was built with the correct metric) rather than from config.

Distance Convention Change — Correct

Knowhere no longer negates IP/cosine distances (faiss now returns true similarity). Sign flip removed from post-processing. Cosine post-scaling math correct for both fp32 and int8 paths. Sentinel handling is metric-appropriate.

GPU_HNSW_SQ Aliasing — Correct

self_key = Type() instead of hardcoded INDEX_GPU_HNSW — correctly handles both base (INDEX_GPU_HNSW) and SQ subclass (INDEX_GPU_HNSW_SQ) binary set keys. New test exercises the exact regression (SQ deserialization would fail without this fix).

New Recall Parity Tests — Well-Designed

gpu_recall >= cpu_recall - 0.05f is a better invariant than absolute floors — correctly handles low-recall-at-small-ef where both CPU and GPU are equally bad. New tests:

• SQ alias regression test
• IP/cosine distance sign test (recomputes true similarity, compares to reported distance within 1e-2 relative error)
• k-valid-results test (no sentinel leakage)

Previously-Flagged Items — All Addressed

• C1 (lazy-upload metric detection): FIXED ✅
• D1 (stale "+128 shift" registration comments): FIXED ✅
• I7 (cudaFuncSetAttribute caching): Fixed (but has TOCTOU race — see faiss PR comment)
• M4 (ScratchSlot destructor): Fixed ✅
• I8 (k<=ef): Fixed ✅

No new critical issues. The C1 fix removes the last blocking item.

[premal]

Code Review — Delta 4 (I1 fix — index-derived metric in post-processing + regression test)

✅ I1 (Important) — FULLY FIXED

The post-processing in Search() now derives cosine detection from gpu_snapshot->isCosine() (the GPU index's upload-time metric) instead of hnsw_cfg.metric_type.value() (the search config). Both post-processing paths — int8 DP4A and fp32 — use the same index-derived is_cosine variable. There are zero remaining config-based metric detection paths in Search().

With C1 fixed in delta 3 (lazy-upload now calls UploadCpuIndexToGpuLocked()) and I1 fixed in delta 4 (post-processing uses isCosine()), the entire metric detection pipeline — upload → search → post-processing — now consistently derives the metric from the GPU index's upload-time state, never from the search config.

Regression Test — Well-Designed

"Test GPU HNSW Cosine Metric Omitted From Search Config" builds a cosine index, searches with metric_type=COSINE (baseline), then searches with metric_type erased from config, and asserts identical ids and distances within 1e-5 epsilon. This directly exercises the I1 scenario and would fail under the old code.

Minor: the test only covers the fp32 path, not int8 DP4A — but both use the same is_cosine variable, so indirect coverage exists.

No New Critical or Important Issues

No merge blockers remain on the knowhere side. All previously-flagged items from our review series are now resolved:

• C1 (lazy-upload metric detection): Fixed ✅
• I1 (post-processing metric detection): Fixed ✅
• I7 (cudaFuncSetAttribute caching TOCTOU): Fixed (mutex) ✅
• D1 (stale "+128 shift" comments): Fixed ✅
• I8 (k<=ef validation): Fixed ✅

[premal]

Code Review — Delta 5 (reset() fix + CMake GPU source documentation)

✅ M1 from Delta 4 — FIXED

reset() now clears is_cosine_ and use_ip_ (mirrors the faiss-side fix). Clean.

CMake Comment Block — Good Documentation

The new comment block in the faiss CMakeLists explains the GPU source compilation architecture clearly:

• Knowhere assembles faiss manually (never add_subdirectory()s the vendored faiss tree)
• faiss_gpu_hnsw OBJECT library is the only place GPU sources are compiled
• The two GPU build paths (upstream faiss CMakeLists vs knowhere's manual list) are mutually exclusive — no duplicate compilation or ODR hazard
• Guidance: "add GPU sources to the list below, NOT by pulling in the vendored gpu/CMakeLists.txt"

This is valuable for future maintainers. No issues.

No new issues. All previously-flagged items remain resolved.

[premal]

Delta 6 Review — knowhere PR #2 (filtered search wiring + tests merged in)

The filtered search wiring from PR #8 has been folded into PR #2. Reviewed the delta against our previous PR #8 review.

Previously-Flagged Items — Most Fixed! ✅

• I1 (dtype/metric coverage): ✅ Substantially addressed. Test now covers FP32 L2/IP/COSINE + INT8 L2/IP/COSINE via GENERATE. The int8 DP4A path (the most distinct code path via searchHostInt8) is covered. Minor gap: FP16/BF16/SQ8 filtered variants still missing (they share the same searchHost fp32 path — low priority).
• I2 (CPU-HNSW baseline + threshold): ✅ Fully addressed. Test now returns {gpu_recall, cpu_recall} and asserts gpu_recall >= cpu_recall - 0.05f (relative parity gate). The old fixed 0.85 brute-force threshold is gone. Ratio sweep expanded to include 0.0 (unfiltered fast path) and 0.3/0.7 (alpha-gate divergence range).
• I4 (ensure_filter device context): ✅ Fixed. cudaSetDevice(device) now present.
• M1 (test comment): ✅ Fixed. Now says "Randomly generated training data" with correct L2 self-distance explanation.
• M3 (edge cases): ✅ Fixed. New [gpu_hnsw_filtered_edge] test covers all-filtered and k > live_count.
• M4 (disable_fallback_brute_force): Still not tested. Carried forward.

Vendoring ✅

All faiss PR #7 additions present in knowhere tree. GpuHnswBruteForce.cuh is byte-identical (162 lines, zero diff). All other files' PR7 additions confirmed present.

New Issues (all minor)

• N1: FP16/BF16/SQ8 filtered recall not tested (shares fp32 code path, low priority)
• N2: disable_fallback_brute_force still untested (carried M4)
• N3: k > live_count edge case doesn't assert exact sentinel count (could mask recall issue)

Verdict

Approved. Great progress on test coverage — the CPU-HNSW parity gate with 0.05 tolerance and the expanded ratio sweep (including 0.0/0.3/0.7) are exactly what was needed. Ready for GPU runtime validation.

Full review at $M/reports/review-knowhere-pr2-delta6.md