feat: wire up GPU_HNSW / GPU_HNSW_SQ index types#6
feat: wire up GPU_HNSW / GPU_HNSW_SQ index types#6devin-ai-integration[bot] wants to merge 45 commits into
Conversation
Adds the GPU_HNSW (and GPU_HNSW_SQ) index type end to end: index-param checker, index-param plumbing, proxy/querynode wiring, the override_index_type swap across all hot-load paths (LoadSegments / ReopenSegments / executor getLoadInfo), GPU Dockerfile deps, KNOWHERE_VERSION pin, design doc, and integration/unit/e2e tests. Feature-only, rebased onto current master (post milvus-io sync). The GPU_TIMING profiling scaffolding is excluded (debug-only). Signed-off-by: Devin AI <devin-ai-integration[bot]@users.noreply.github.com>
🤖 Devin AI EngineerI'll be helping with this pull request! Here's what you should know: ✅ I will automatically:
Note: I can only respond to comments from users who have write access to this repository. ⚙️ Control Options:
|
cmake 4.3.4 rejects the literal \u2014 (intended em-dash) in three install-warning messages as an invalid escape sequence. Use -- instead. Introduced by the knowhere v3.0.5 upgrade (upstream 1065ab0). Signed-off-by: Devin AI <devin-ai-integration[bot]@users.noreply.github.com>
- gpuIndexRequiresGpu now returns true for GPU_HNSW / GPU_HNSW_SQ so the segment loader runs the pre-load VRAM check and the collection is flagged for GPU task scheduling (previously fell through default->false) - LoadIndex (hot index-update path) applies AppendPrepareLoadParams so the override_index_type swap is present even if the handler-level defense is bypassed - lb_policy.go: revert no-op split-return refactor to plain returns - comment the intentional idempotent double AppendPrepareLoadParams in Load - design doc: add MEP metadata header per CONTRIBUTING.md KNOWHERE_VERSION intentionally still points to the 6si/knowhere gpu-hnsw-faiss branch while the feature is unmerged; to be repinned to a tagged release before merge. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…n record The overflow candidate queue was never functional (overflow_insert had no call site; d_overflow_count only zeroed+read). Removed from the kernel in 6si/faiss#1 and 6si/knowhere#2. Update the design doc to describe a plain parallel beam search and add a decision record explaining the removal. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Resolve review findings on the GPU_HNSW PR: - Register newGpuHnswChecker() for GPU_HNSW / GPU_HNSW_SQ and route those two types to it ahead of the unified vecIndexChecker short-circuit, so the M/efConstruction range validation actually runs (knowhere's ValidateIndexParams is a no-op for these types). Previously the checker was defined but never registered and thus unreachable. - Fix the checker unit tests that silently returned (reported PASS) when GetChecker yielded nil: they now t.Skip honestly on non-GPU builds, and a pure-Go routing test asserts GetChecker resolves to *gpuHnswChecker. - Design doc: remove the stale internal/datacoord/meta.go compaction claim (normalizePositionTimestamp is upstream and not in this PR), add the filtered-search (delete/TTL/partition bitset) limitation to Failure Modes, and record the GPU_HNSW_SQ-alias and filtered-search decisions in the decision record. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
GPU_HNSW is registered in knowhere only for FLOAT32 (GPU_ANN_FLOAT_INDEX = FLOAT32 | GPU) and INT8, so CheckValidDataType correctly rejects Float16Vector and BFloat16Vector. The test asserted both were valid (a stale claim of fp16/bf16 support); align expectations with the actual/intended fp32+int8 scope. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Knowhere now advertises FP16/BF16 for GPU_HNSW and GPU_HNSW_SQ (feature::FP16|GPU, feature::BF16|GPU), so the checker's CheckValidDataType auto-accepts them via the cgo feature flags — no Go logic change needed. - gpu_hnsw_checker_test.go: revert the earlier expectation that Float16/ BFloat16 vectors are invalid; both are now expected valid. - design docs: document native 2-byte fp16/bf16 device storage with per-element up-conversion to fp32 in the kernel, and add the native FP16/BF16 decision record (context, options, compile-only verification status). Compile/checker-verified only; GPU-runtime validation pending capacity. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…gments ReopenSegments ran requestResource() before applying the knowhere override_index_type, so the memory-budget check saw the CPU index type (HNSW, memoryCost=file_size) instead of the GPU type (GPU_HNSW, memoryCost=0), and could spuriously reject a GPU reopen with an insufficient-memory error. Move the AppendPrepareLoadParams loop before requestResource, matching the ordering already used in Load() and LoadIndex(). AppendPrepareLoadParams is idempotent, so segment.Reopen still receives the overridden params. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
GPU_HNSW frees its CPU copy after uploading the graph/vectors to VRAM, so its retained memoryCost is ~0 and the previous non-mmap reservation of 2*memoryCost reserved zero bytes even though Deserialize() transiently materializes the CPU HNSW index plus decode/graph staging buffers before the upload. Concurrent loads could OOM the query node. Knowhere's Resource now exposes a phase-separated maxMemoryCost (peak transient host memory during load). When non-zero, use it for the non-mmap max_memory_cost reservation; final_memory_cost stays res.memoryCost. Indexes that do not set the field (maxMemoryCost==0) keep the existing 2*memoryCost heuristic, and the disk/mmap path is unchanged. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Codex P1 #3 — reserve GPU_HNSW peak host memory (
|
KNOWHERE_VERSION pointed at the mutable branch gpu-hnsw-faiss, so the knowhere build could change with no milvus PR. Pin to the current branch-HEAD SHA 22fd143b360a0b37bc89a2128c496b95abfb3914 for reproducibility. Must be repointed to a tagged 6si/knowhere release before merging to production. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Codex re-review should-fix #1 — pin KNOWHERE_VERSION to a commit SHA (
|
…n test) Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…-upload nit) Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
GPU_HNSW uploads the CPU index + fp32 staging to host RAM (~2x file size) during Deserialize and frees it once vectors reach VRAM, so its *resident* host cost is ~0. Under tiered eviction the caching layer (MCL) tracks only resident cost, and segment_loader skipped charging maxMemoryCost to the loading estimate, so GPU_HNSW segments were admitted against a ~0-cost estimate. With many parallel loads the accumulated transient peaks OOMKilled the querynode (regression introduced when tiered eviction was enabled). Charge a GPU index's maxMemoryCost (the transient host peak) to the loading estimate even under tiered eviction, via indexLoadingMemoryContribution(). The reservation flows through TryReserveLoadingResourceWithTimeout and is released in freeRequestResource once the load completes, so it throttles concurrent loads without inflating steady-state resident cost. Non-GPU indexes under tiered eviction are unchanged (contribution 0); non-tiered behavior is unchanged. Adds TestIndexLoadingMemoryContribution covering all four cases. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…oad fix) Picks up 6si/knowhere#2 6906c993, which runs the eager GPU upload on the DeserializeFromFile (mmap) load path that Milvus uses for sealed segments — fixing the querynode host-RAM OOM where the CPU-built HNSW was retained in host memory instead of being uploaded to VRAM and freed. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…re pin The FieldIndexInfo frequently omits NumRows (it carries index metadata, not segment stats), so the C load info was built with NumRows=0. Knowhere's load-resource estimator multiplies rows*dim to size the transient host peak of a GPU_HNSW upload; with rows=0 every row-based term vanished and the estimate collapsed to ~file_size, letting the loader admit far too many concurrent uploads and OOM the querynode host cgroup before any upload completed. Fall back to loadInfo.NumOfRows (the sealed-segment row count, equal to the indexed vector count) when the index info's row count is absent. Bump KNOWHERE_VERSION to 4d3d97d8, which sizes the estimate to the fp32 resident graph and restores mmap eligibility for GPU_HNSW/GPU_HNSW_SQ. issue: #6 Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…estimate) Picks up the knowhere fix that sizes the GPU_HNSW load transient at ~2x file (compact int8) instead of the fp32 ~20 GiB/seg overestimate, so the querynode soft-OOM admission guard reflects the true compact footprint. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…rtions) Update pin from 7d8a448a to db139ada which fixes the load-resource test in test_gpu_search.cc — the test assertions were written for a superseded fp32-expansion estimate and would fail on any CUDA-enabled build. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com> Signed-off-by: Premal Shah <premal@6sense.com>
issue: #6 Repoint KNOWHERE_VERSION 5b8baf98 -> 8fbaf1f4 (6si/knowhere@gpu-hnsw-faiss). Pulls in the bitonic_sort_staging shared-memory race fix (real root cause of the v82 fp16 P1 OOB) on top of the node-id OOB guards and per-dtype tests. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Code Review — Delta 2The only change in this delta is bumping All previously flagged items remain open:
No new issues found. |
…, auto-clamp) Repoint KNOWHERE_VERSION 8fbaf1f4 -> fcdcacf2 (6si/knowhere@gpu-hnsw-faiss). fcdcacf2 corrects the IP/COSINE public distance sign at the faiss copy-out boundary (and drops knowhere's compensating re-negation), fixes GPU_HNSW_SQ deserialize aliasing, detects the metric from the FAISS index on lazy upload, validates ef/search_width and auto-clamps ef>=k in the GPU kernel, and adds focused GPU regression tests. No milvus source change beyond the pin. Revert: repoint pin to v80 knowhere 5a5c50285e437424e1a042d55d56134e072b03d7 or redeploy image gpu-hnsw-v80. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…rity test) Repoints the 6si/knowhere fork pin to d022e898, which fixes the [gpu_hnsw_compare] 'Varying ef comparison' test to assert GPU-vs-CPU recall parity instead of an invalid absolute 0.5 floor at low ef. No runtime code change. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…v85) Sync upstream master (41eacb5) into the GPU HNSW branch and re-port the GPU load/VRAM-admission logic onto master's restructured segment loader. Resolutions: - thirdparty/knowhere/CMakeLists.txt: keep 6si/knowhere pin d022e898 (v84 gpu_hnsw_compare parity fix), not upstream zilliztech 49507ef3. - querynodev2/segments/segment_loader.go: apply override_index_type before resource estimation in Load() and master's new ReopenSegments() path; charge GPU_HNSW/GPU_HNSW_SQ transient host upload peak even under tiered eviction via new indexLoadingMemoryContribution helper; gpuIndexRequiresGpu recognizes GPU_HNSW/GPU_HNSW_SQ. - querynodev2/services.go: apply load-stage overrides at LoadSegments entry (defense-in-depth) using mlog/merr. - querynodev2/segments/segment.go: NumRows fallback to segment row count so GPU upload host-peak estimate is not collapsed to file size. Dead getFieldType/LoadIndex paths removed by master milvus-io#51053/milvus-io#51323 are gone; index-update now flows through ReopenSegments, which carries the override. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Code Review — Delta 3The only change is bumping All previously flagged items remain open:
No new issues found. |
… e2e - Bump the knowhere pin to 6si/knowhere@898a93f2, which fixes the GPU HNSW query-side metric detection (derive cosine/IP from the built GPU index, not the search config) and re-vendors the faiss GPU tree with the thread-safe cudaFuncSetAttribute smem configuration (TOCTOU fix) and entry_point bounds validation. The pin stays an immutable fork SHA; the existing comment already flags repointing to a tagged 6si/knowhere release before a production merge. - Replace fixed time.sleep() guesses in test_gpu_hnsw_e2e.py with condition waits: wait_for_loading_complete() for load/reload, a release-state poll, and a search-based poll for the hot-loaded sealed segment. Removes flakiness and avoids over-waiting. Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Code Review — Delta 4 (time.sleep replaced with polling + SHA bump)✅ M6 (Minor) — FIXED
The hot-segment test now polls until the hot-inserted vector self-searches to its own ID (
KNOWHERE_VERSION bumped to
|
…st M<=512); bump knowhere pin Two create-time param-checker bugs surfaced in PR review of gpuHnswChecker: - M and efConstruction were mandatory: CheckIntByRange returns false for an absent key, which was turned into errOutOfRange, so a normal create_index that omits them (relying on knowhere defaults, as CPU HNSW allows) was refused. Now range-checked only when supplied. - Max M contradiction: the range check advertised M up to HNSWMaxM=2048, but a separate staging guard rejected M>512 (next_pow2(2*M) must fit one CUDA block). Reconciled to the honest GPU bound: gpuHnswMaxM=512, enforced in the range check itself; M>512 builds but is never GPU-searchable, so it is rejected at create. CPU HNSW keeps its 2048 max. Updated the Go checker unit test (M=512 boundary, M=513 rejected, optional M/efConstruction) and the python index specs (idx_gpu_hnsw / idx_gpu_hnsw_sq: max boundary 2048 -> 512). Bumped KNOWHERE_VERSION to da32be6e (faiss reset() metric-flag clear, re-vendored). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…d range Emit the same 'param <key> (<v>) should be in range [<min>, <max>]' wording knowhere uses for CPU HNSW (knowhere config validation is a no-op for GPU_HNSW, so this Go checker is the authoritative validator), and use gpuHnswMinM=2 so the enforced M range [2, 512] matches the message and the python index specs exactly. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
… path - Add a prominent top-of-doc callout that GPU_HNSW/GPU_HNSW_SQ does not apply bitsets (deletes/TTL/partition), so it is append-mostly/immutable only. - Add a Failure-Modes row documenting that vector-field output is served from raw field data (not the index's GetVectorByIds): IndexHasRawData=false makes segcore eager-load raw field data and route retrieval through get_raw_data. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Code Review — Delta 5 (GPU HNSW M/efConstruction param validation + tests + doc updates)SummaryNo critical or important issues. This delta is clean and well-implemented. ✅ GPU M Validation — Correct
✅ Optional Parameter Semantics — CorrectM and efConstruction are only validated when present ( ✅
|
| Finding | Status |
|---|---|
| I1 (maxMemoryCost for VRAM admission) | Unchanged (non-blocking) |
| I2 (KNOWHERE_VERSION fork SHA) | Bumped to da32be6e, still 6si fork |
| M6 (time.sleep in E2E tests) | Fixed in delta 4, no regression |
No merge blockers.
Live e2e (v87) showed: M=513 is caught by the Go GPU checker (err 1100, 'param M (513) should be in range [2, 512]'), while M=-1 is caught earlier by knowhere's stock HNSW JSON validator (err 1100, 'Out of range in json: param M (-1) should be in range [2, 2048]'). Both reject correctly; only the layer/message differs. Update the negative-M expected msg to [2, 2048] (was [2, 512], which no longer substring-matches) and set err_code to 1100. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
CPU HNSW allows M=1 (HNSWMinM=1); the GPU kernel requires M>=2. The comment said 'matches the CPU HNSW minimum' — corrected to 'stricter than', mirroring the gpuHnswMaxM comment. Comment-only; no behavior change. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
|
Delta-5 minor (gpuHnswMinM comment): fixed in 5061f5e — the comment now says the GPU M>=2 minimum is stricter than the CPU HNSW minimum (HNSWMinM=1), mirroring the gpuHnswMaxM comment. Comment-only, no behavior change. |
…earch Specifies GPU-side bitset filtering for GPU_HNSW/GPU_HNSW_SQ to reach CPU HNSW parity: two-tier beam (valid results vs. invalid frontier), the accumulated_alpha/kAlpha admission gate, and brute-force fallback at high filter ratios (kHnswSearchKnnBFFilterThreshold=0.93, kHnswSearchBFTopkThreshold=0.5). Scope is faiss kernel + knowhere plumbing; Milvus delete/L0/bitset pipeline is unchanged (docs/tests only). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
- I1: concrete two-tier smem layout (ef_valid=ef unchanged; validity packed into is_expanded; separate ef_inv frontier; max_ef recomputed at the real 24 B/slot base, not 12) - I2: device-side needs_bf[] short-query fallback (atomic append + one batched BF launch on the same stream) instead of a mid-search host round-trip - I3: batch-level alpha gate applied in the post-bitonic-sort serial section (reviewer Option 3), preserving CPU rate-limiting effect deterministically - M1-M6: BF shares result scratch; stream-ordering/lifetime invariant; d_bitset realloc on N change; search_width/staging unaffected; disable_fallback_brute_force threading; new BF kernel file added to the knowhere GPU CMake source list Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…e-verified) - N1: cross-beam merge/eviction/parent-selection decision tree (5.3.1) - filtered->alpha gate->invalid frontier (dist < valid worst); parent selection picks globally-closest unexpanded across both beams (pop_based_on_distance analog); termination mirrors CPU has_next() - needs_bf_count zeroed via cudaMemsetAsync on slot stream; accumulated_alpha carried in meta[] init 1.0f (matches CPU non-BF init) - Rejected the '12 B/slot' correction after re-reading source: six ef-length arrays (result_* + merged_*) and a /24 divisor; merge is not in-place. Base stays 24 B/ef; max_ef ~= (smem-overhead)/36 with the invalid frontier Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…r GPU-side deletes The GPU image builds knowhere via FetchContent from KNOWHERE_VERSION. The prior pin (da32be6e) predates GPU-side filtered search, so the packaged libknowhere.so rejected non-empty bitsets and deletes on sealed GPU_HNSW segments returned invalid_args. Bump to 45e2ba26 (6si/knowhere@gpu-hnsw-filtered-search) which implements the two-tier beam + alpha gate + BF fallback validated by the knowhere GPU unit gate (G-ID, filtered recall FP32/INT8, edge, dtype matrix) and clean memcheck/initcheck/racecheck. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
enhance(gpu_hnsw): bump knowhere pin to filtered-search (45e2ba26) for GPU-side deletes
Delta 6 Review — milvus PR #6 (filtered search docs + tests)Previously-Flagged Items
New Issues — Two Merge BlockersN1: Design docs NOT flipped, contradicting the version bump. The MEP §8 explicitly instructs flipping the "not supported" callouts in Fix: Flip the doc callouts as the MEP instructs, or revert the version bump. N2: No delete-then-query integration tests. The MEP §8 instructs flipping the Fix: Add delete-then-query test cases that verify filtered rows are absent from results and non-deleted rows are present. Root CauseBoth N1 and N2 stem from bumping KNOWHERE_VERSION without applying the Milvus-side doc/test changes the MEP requires. Either flip the docs + add the tests, or revert the version bump until the Milvus-side changes are ready. Positive
Full review at $M/reports/review-milvus-pr6-delta6.md |
…e-then-query integration tests Addresses the milvus#6 review blockers (N1/N2): the KNOWHERE_VERSION bump now carries GPU-side filtered search, so the docs must no longer say GPU_HNSW rejects bitsets and there must be integration coverage for delete-then-query. - docs: flip the 'not supported / returns invalid_args' callouts in 20260619-gpu-hnsw.md (key-limitation banner + Failure Modes row) to 'supported (CPU parity; BF fallback)'; mark the filtered-search decision in 20260711-gpu-hnsw-ocq-removal.md as superseded by 20260718; set 20260718 status to Implemented. - tests: add TestGpuHnswFilteredSearch / TestGpuHnswSQFilteredSearch — insert + flush (seal) + GPU index + delete a PK range, then assert deleted rows are absent from search and query, live rows remain, and no invalid_args, across L2/IP/COSINE with Strong consistency. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
|
Addressed the two merge blockers in N1 — docs flipped. N2 — delete-then-query integration tests added. The minor items (faiss N1 codegen-identical comment, knowhere M4/N2 disable_bf/N1 dtype UT, N3 sentinel-count) are non-blocking and left as noted. |
…istance Repoints KNOWHERE_VERSION 45e2ba26 -> 407b6095 (6si/knowhere@gpu-hnsw-faiss), which re-vendors faiss PR #8: layer-0 neighbor expansion is now warp-per-candidate cooperative distance (coalesced dataset loads) instead of thread-per-candidate. Kernel-only measurements on the L40S build box (counter deltas transfer to the Blackwell querynodes): sectors/request 12.1 -> 3.10 (3.9x), throughput 1.57x int8-DP4A / 1.85x fp32, recall identical to baseline (7/7, max delta -0.0005), compute-sanitizer clean. No milvus-side code or API change; pin bump only. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
… fix Pins knowhere to 2b52090d (6si/knowhere#10), which re-vendors the faiss GPU HNSW fix bounding the grow-only visited bitmap via nq-chunking. Supersedes the H1 pin in #12 for the v91 cluster build/validation; re-pin to the merged knowhere SHA before merge. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Clarify the pin comment now covers the nq-chunked visited-bitmap OOM fix in addition to H1. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…mmit Repoint from the pre-merge feature-branch SHA 2b52090d to f7126691, the merge commit of 6si/knowhere#10 on gpu-hnsw-faiss, so the pin references the permanent integration branch rather than a feature branch. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
enhance(gpu_hnsw): bump knowhere pin to nq-chunked visited-bitmap OOM fix
Summary
Wire up the
GPU_HNSW/GPU_HNSW_SQindex types in Milvus: load-timeoverride_index_typeswap (HNSW/HNSW_SQ → GPU_HNSW) across the load, hot-load, and querycoord paths; GPU resource accounting/routing; and create-time param validation. Depends on 6si/knowheregpu-hnsw-faiss(pinned viaKNOWHERE_VERSION; repin to a tagged release before production merge).Core wiring:
segment_loader.go/services.go/querycoordv2/task/executor.go— apply theoverride_index_typeswap before resource estimation and index load.gpuIndexRequiresGpureturns true forGPU_HNSW/GPU_HNSW_SQ→ pre-load VRAM check runs and the collection is flagged for GPU scheduling.Data types — native FP32/FP16/BF16/INT8: GPU_HNSW/GPU_HNSW_SQ now accept fp16/bf16 in addition to fp32/int8. This is entirely feature-flag driven — no Go logic change. Knowhere advertises
feature::FP16|GPU/feature::BF16|GPUfor both types (see 6si/knowhere#2), and Milvus reads those bits via cgo:The only change is reverting the earlier
gpu_hnsw_checker_test.goexpectation (fp16/bf16 were temporarily marked invalid when knowhere was fp32/int8-only); both are now expected valid. The python param specs (idx_gpu_hnsw.py/idx_gpu_hnsw_sq.py) already listFLOAT16_VECTOR/BFLOAT16_VECTOR.Param checker (review fix #4):
newGpuHnswChecker()was defined but never registered, andGetChecker()short-circuits any vec-index to the unifiedvecIndexChecker(whose knowhereValidateIndexParamsis a no-op for GPU_HNSW), so M/efConstruction validation never ran. Now:registerIndexChecker()registersnewGpuHnswChecker()for both types. The 3 checker tests that silentlyreturned (reported PASS) whenGetCheckeryielded nil nowt.Skiphonestly on non-GPU builds, plus a pure-Go routing test assertsGetCheckerresolves to*gpuHnswChecker.Design doc (review fixes #3, #5 + native fp16/bf16):
internal/datacoord/meta.goclaim —normalizePositionTimestampis upstream andmeta.gois not in this PR; GPU_HNSW just inherits the upstream compaction behavior.BitsetViewon the GPU (filtered rows excluded from results but kept as graph waypoints; brute-force fallback at high filter ratios), validated at CPU-HNSW parity. The design docs (20260619-gpu-hnsw.mdkey-limitation banner + Failure Modes row) are flipped from "not supported / returns invalid_args" to "supported", the20260711-gpu-hnsw-ocq-removal.mdreject decision is marked superseded by20260718-gpu-hnsw-filtered-search.md, and delete-then-query integration tests were added (TestGpuHnswFilteredSearch/TestGpuHnswSQFilteredSearchintests/python_client/testcases/indexes/test_gpu_hnsw{,_sq}.py): insert → flush(seal) → GPU index → delete a PK range → assert deleted rows absent from search+query, live rows present, noinvalid_args, across L2/IP/COSINE.docs/design-docs/design_docs/20260711-gpu-hnsw-ocq-removal.md.GPU_HNSW_SQ (review fix #2): the type Milvus exposes now has a real Knowhere backend (alias of GPU_HNSW; see 6si/knowhere#2) instead of hitting no impl.
Verification: not build-verified locally (no Go/CUDA build in this session); CI / the GPU build confirm compilation. The
Test_gpuHnswChecker_CheckValidDataTypefp16/bf16 assertions require the linked core to include the updated knowhere registrations (the next GPU build). GPU-runtime validation of native fp16/bf16 is covered by the GPU unit/recall gates and the sealed-segment cluster e2e (see Performance below).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):Throughput figures are from the
gpu-hnsw-v89image (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, 6si/knowhere#2.
Link to Devin session: https://6sense.devinenterprise.com/sessions/d39aba56b8a3467cbbf231ab631a06ed
Link to Devin session: https://6sense.devinenterprise.com/sessions/55dab0bd33c346df99b5818b30059342
Requested by: @premal