Skip to content

feat: wire up GPU_HNSW / GPU_HNSW_SQ index types#6

Open
devin-ai-integration[bot] wants to merge 45 commits into
masterfrom
gpu-hnsw-faiss
Open

feat: wire up GPU_HNSW / GPU_HNSW_SQ index types#6
devin-ai-integration[bot] wants to merge 45 commits into
masterfrom
gpu-hnsw-faiss

Conversation

@devin-ai-integration

@devin-ai-integration devin-ai-integration Bot commented Jul 9, 2026

Copy link
Copy Markdown

Summary

Wire up the GPU_HNSW / GPU_HNSW_SQ index types in Milvus: load-time override_index_type swap (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/knowhere gpu-hnsw-faiss (pinned via KNOWHERE_VERSION; repin to a tagged release before production merge).

Core wiring:

  • segment_loader.go / services.go / querycoordv2/task/executor.go — apply the override_index_type swap before resource estimation and index load.
  • gpuIndexRequiresGpu returns true for GPU_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|GPU for both types (see 6si/knowhere#2), and Milvus reads those bits via cgo:

// vecindexmgr: Float16Flag=1<<2, BFloat16Flag=1<<3 match knowhere feature::FP16/BF16
C.GetIndexFeatures(...)  // populates mgr.features[GPU_HNSW]
// -> IsDataTypeSupport(GPU_HNSW, Float16Vector/BFloat16Vector) now true
//    -> checker.CheckValidDataType accepts them

The only change is reverting the earlier gpu_hnsw_checker_test.go expectation (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 list FLOAT16_VECTOR/BFLOAT16_VECTOR.

Param checker (review fix #4): newGpuHnswChecker() was defined but never registered, and GetChecker() short-circuits any vec-index to the unified vecIndexChecker (whose knowhere ValidateIndexParams is a no-op for GPU_HNSW), so M/efConstruction validation never ran. Now:

// GetChecker: prefer the dedicated checker for these two types
switch indexType {
case IndexGpuHnsw, IndexGpuHnswSQ:
    if c, ok := mgr.checkers[indexType]; ok { return c, nil }
}
// ...then the unified vec-index short-circuit

registerIndexChecker() registers newGpuHnswChecker() for both types. The 3 checker tests that silently returned (reported PASS) when GetChecker yielded nil now t.Skip honestly on non-GPU builds, plus a pure-Go routing test asserts GetChecker resolves to *gpuHnswChecker.

Design doc (review fixes #3, #5 + native fp16/bf16):

  • Removed the stale internal/datacoord/meta.go claim — normalizePositionTimestamp is upstream and meta.go is not in this PR; GPU_HNSW just inherits the upstream compaction behavior.
  • Filtered search — now supported (CPU parity). GPU_HNSW/GPU_HNSW_SQ apply the delete/TTL/partition BitsetView on 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.md key-limitation banner + Failure Modes row) are flipped from "not supported / returns invalid_args" to "supported", the 20260711-gpu-hnsw-ocq-removal.md reject decision is marked superseded by 20260718-gpu-hnsw-filtered-search.md, and delete-then-query integration tests were added (TestGpuHnswFilteredSearch / TestGpuHnswSQFilteredSearch in tests/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, no invalid_args, across L2/IP/COSINE.
  • Documented native 2-byte fp16/bf16 device storage (up-converted to fp32 in the kernel) and recorded the GPU_HNSW_SQ-alias, filtered-search, checker, and native-FP16/BF16 decisions in 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_CheckValidDataType fp16/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):

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, 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

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-integration

Copy link
Copy Markdown
Author

🤖 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

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>
devin-ai-integration[bot]

This comment was marked as resolved.

devin-ai-integration[bot]

This comment was marked as resolved.

devin-ai-integration[bot]

This comment was marked as resolved.

devin-ai-integration[bot]

This comment was marked as resolved.

devin-ai-integration[bot]

This comment was marked as resolved.

devin-ai-integration[bot]

This comment was marked as resolved.

premal and others added 2 commits July 11, 2026 20:08
- 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>
devin-ai-integration[bot]

This comment was marked as resolved.

premal and others added 4 commits July 11, 2026 22:06
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>
@devin-ai-integration

Copy link
Copy Markdown
Author

Codex P1 #3 — reserve GPU_HNSW peak host memory (5e13c42763)

GPU_HNSW frees its CPU copy after uploading graph/vectors to VRAM, so its retained memoryCost is ~0. The previous non-mmap reservation max_memory_cost = 2 * memoryCost therefore reserved zero bytes, even though Deserialize() transiently materializes the CPU HNSW index + decode/graph staging 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 — 6si/knowhere#2 08ced662). VectorIndexLoadResource in IndexFactory.cpp:

} else if (res.maxMemoryCost > 0) {          // GPU_HNSW: peak != retained
    request.max_disk_cost = 0;
    request.max_memory_cost = res.maxMemoryCost;
} else {                                       // unchanged for all other indexes
    request.max_disk_cost = 0;
    request.max_memory_cost = 2 * res.memoryCost;
}

final_memory_cost stays res.memoryCost; disk/mmap path unchanged; indexes that don't set the field keep the 2*memoryCost heuristic. This composes with the earlier ReopenSegments override-ordering fix (b6bd313d) so the GPU_HNSW estimator — not the CPU HNSW one — is what feeds the reservation.

This is a C++ recompile of segcore (needs to build against knowhere 08ced662); v67 build queued. Assertion coverage for the accounting is in the knowhere [gpu_hnsw_p1] "Load-resource estimate" section.

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>
@devin-ai-integration

Copy link
Copy Markdown
Author

Codex re-review should-fix #1 — pin KNOWHERE_VERSION to a commit SHA (0c81fc3a8b)

internal/core/thirdparty/knowhere/CMakeLists.txt pointed KNOWHERE_VERSION at the mutable branch gpu-hnsw-faiss, so the knowhere build could change with no milvus PR. Now pinned to the immutable branch-HEAD SHA 22fd143b360a0b37bc89a2128c496b95abfb3914 (FetchContent GIT_TAG = full SHA, full clone). Still a pre-merge action item: repoint to a tagged 6si/knowhere release once one exists (comment left in the file). v68 build will confirm FetchContent resolves the SHA cleanly.

premal and others added 8 commits July 12, 2026 02:55
…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>
@premal

premal commented Jul 17, 2026

Copy link
Copy Markdown

Code Review — Delta 2

The only change in this delta is bumping KNOWHERE_VERSION to 8fbaf1f4746c54bb71a90865e1d45f081426533b, which picks up the knowhere bitonic sort race fix + merge-path off-by-one fix + defensive guards. No new milvus code changes.

All previously flagged items remain open:

  • [I1] fieldGpuMemorySize charged MaxMemoryCost (host-RAM peak) for VRAM admission
  • [I2] KNOWHERE_VERSION still on 6si fork, comment-only repoint action item
  • [M6] E2E tests use time.sleep() instead of proper wait conditions

No new issues found.

premal and others added 3 commits July 17, 2026 14:56
…, 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>
@premal

premal commented Jul 18, 2026

Copy link
Copy Markdown

Code Review — Delta 3

The only change is bumping KNOWHERE_VERSION to d022e89825c499d43822c90041ebb9d4908b6a25 (picks up knowhere C1 fix + distance convention change + SQ alias fix). Also reverted the knowhere fork URL/git repo back to zilliztech — the version is now pinned to a 6si fork commit SHA but the repo URL points to the upstream. No new milvus code changes.

All previously flagged items remain open:

  • [I1] fieldGpuMemorySize charged MaxMemoryCost (host-RAM peak) for VRAM admission
  • [I2] KNOWHERE_VERSION still on 6si fork SHA, comment-only repoint action item
  • [M6] E2E tests use time.sleep() instead of proper wait conditions

No new issues found.

@devin-ai-integration devin-ai-integration Bot left a comment

Copy link
Copy Markdown
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Devin Review found 2 new potential issues.

View 3 additional findings in Devin Review.

Open in Devin Review

Comment thread internal/util/indexparamcheck/gpu_hnsw_checker.go Outdated
Comment thread internal/util/indexparamcheck/gpu_hnsw_checker.go Outdated
… 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>
@premal

premal commented Jul 18, 2026

Copy link
Copy Markdown

Code Review — Delta 4 (time.sleep replaced with polling + SHA bump)

✅ M6 (Minor) — FIXED

time.sleep() calls in E2E tests replaced with proper polling:

  • wait_load_complete() — delegates to SDK's wait_for_loading_complete
  • wait_until(predicate, timeout, interval, desc) — polls an observable condition with configurable timeout

The hot-segment test now polls until the hot-inserted vector self-searches to its own ID (_hot_segment_searchable), instead of guessing time.sleep(5). This is the correct observable condition — race-free and doesn't over-wait.

wait_until is well-designed: configurable timeout/interval, includes last exception in timeout error, good description for diagnostics.

KNOWHERE_VERSION bumped to 898a93f2

Picks up knowhere isCosine() API + I1 fix. Still on 6si fork SHA.

Remaining Open Items (non-blocking)

  • [I1] fieldGpuMemorySize charged MaxMemoryCost (host-RAM peak) for VRAM admission — verify this bounds VRAM for GPU_HNSW
  • [I2] KNOWHERE_VERSION on 6si fork SHA — repoint to tagged release before production merge

No new issues found.

premal and others added 3 commits July 18, 2026 14:54
…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>
@premal

premal commented Jul 18, 2026

Copy link
Copy Markdown

Code Review — Delta 5 (GPU HNSW M/efConstruction param validation + tests + doc updates)

Summary

No critical or important issues. This delta is clean and well-implemented.

✅ GPU M Validation — Correct

gpuHnswMaxM = gpuHnswMaxStagingCapacity / 2 = 512 is the right bound. The GPU search kernel stages next_pow2(2*M) candidates in one CUDA block (max 1024 threads). M=512 → next_pow2(1024)=1024 ✓, M=513 → next_pow2(1026)=2048 > 1024 ✗. Good simplification from the previous runtime nextPow2() to a compile-time constant.

✅ Optional Parameter Semantics — Correct

M and efConstruction are only validated when present (if _, ok := params[key]; ok), matching CPU HNSW behavior where knowhere fills defaults. Tests verify both-omitted, M-omitted, and ef-omitted all succeed.

errParamShouldBeInRange — Consistent with Knowhere

"param '%s' (%s) should be in range [%d, %d]" mirrors knowhere's CPU HNSW error wording. Python tests verify exact messages.

✅ Test Coverage — Comprehensive

Go tests: boundary (M=512 pass, M=513 fail, M=-1 fail, M=600 fail), optional semantics, non-power-of-2 M, metric types, data types. Python tests: M boundaries, type validation, efConstruction boundaries, SQ-specific params. Good coverage.

✅ Design Doc — Accurate

Filtered-search limitation blockquote correctly describes: no BitView support, invalid_args reject for non-empty bitsets, no CPU fallback (index freed after upload), append-only recommendation. Vector-output safety row correctly explains raw data comes from column data, not the index.

Minor: M1 — gpuHnswMinM Comment Inaccurate

// gpuHnswMinM is the smallest usable HNSW graph degree (a graph needs at least
// 2 neighbors per node); matches the CPU HNSW advertised minimum.
const gpuHnswMinM = 2

CPU HNSW allows M=1 (HNSWMinM = 1 in constraints.go). GPU requires M≥2. The comment should say "stricter than the CPU HNSW minimum" rather than "matches," mirroring the gpuHnswMaxM comment which correctly says "stricter than the CPU HNSW max."

Previous Findings Status

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.

premal and others added 2 commits July 18, 2026 19:08
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>
@devin-ai-integration

Copy link
Copy Markdown
Author

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.

premal and others added 5 commits July 18, 2026 20:08
…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
@premal

premal commented Jul 19, 2026

Copy link
Copy Markdown

Delta 6 Review — milvus PR #6 (filtered search docs + tests)

Previously-Flagged Items

  • M1 (delta5 — gpuHnswMinM comment): ✅ Fixed. Comment now correctly says "stricter than the CPU HNSW minimum (HNSWMinM=1)" instead of the incorrect "matches the CPU HNSW advertised minimum."
  • I1 (fieldGpuMemorySize / MaxMemoryCost): Not addressed — non-blocking, carried forward.
  • I2 (KNOWHERE_VERSION on 6si fork SHA): ✅ Bumped from da32be6e...45e2ba26683ec1a62692d7c55dc47da7648e121f, with thorough comment explaining the SHA carries GPU-side filtered search.

New Issues — Two Merge Blockers

N1: Design docs NOT flipped, contradicting the version bump.

The MEP §8 explicitly instructs flipping the "not supported" callouts in 20260619-gpu-hnsw.md and 20260711-gpu-hnsw-ocq-removal.md to "supported (CPU-parity; BF fallback)." But the docs still say "not supported" and "returns invalid_args." Meanwhile KNOWHERE_VERSION is bumped to a SHA that carries filtered search — so the built binary will accept bitsets while the docs tell users it rejects them. Internally inconsistent.

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 idx_gpu_hnsw* Python specs from expecting invalid_args to expecting "deleted rows absent + non-deleted rows present." There are zero delete-then-query tests in the diff. The filtered search path is untested at the Milvus integration level.

Fix: Add delete-then-query test cases that verify filtered rows are absent from results and non-deleted rows are present.

Root Cause

Both 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

  • M1 fix is clean and correct
  • The MEP doc is the exact approved Rev 3
  • CMake comment for KNOWHERE_VERSION bump is thorough
  • UTC import compatibility fix is good defensive coding

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>
@devin-ai-integration

Copy link
Copy Markdown
Author

Addressed the two merge blockers in f9ed9e0185:

N1 — docs flipped. 20260619-gpu-hnsw.md: the key-limitation banner and the Failure Modes row now read "Filtered search — supported (CPU parity)" (filtered rows excluded from results but kept as graph waypoints; brute-force fallback at high ratios) instead of "not supported / returns invalid_args". 20260711-gpu-hnsw-ocq-removal.md: the filtered-search reject decision (Status/Decision/Follow-up) is marked SUPERSEDED by 20260718-gpu-hnsw-filtered-search.md, whose status is now Implemented. Docs are now consistent with the bumped KNOWHERE_VERSION.

N2 — delete-then-query integration tests added. TestGpuHnswFilteredSearch (test_gpu_hnsw.py) and TestGpuHnswSQFilteredSearch (test_gpu_hnsw_sq.py): insert → flush (seal) → build GPU index → load → delete a PK range → assert (a) deleted ids absent from search, (b) query over the deleted range is empty, (c) live rows still returned, (d) no invalid_args — parametrized over L2/IP/COSINE with Strong consistency.

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.

premal and others added 5 commits July 19, 2026 23:54
…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
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant