feat(gpu_hnsw): GPU_HNSW / GPU_HNSW_SQ on faiss GpuIndexHNSW (int8/DP4A, fp16/bf16, filtered, accurate int8 cosine)#11
Closed
devin-ai-integration[bot] wants to merge 44 commits into
Closed
Conversation
GPU_HNSW (and GPU_HNSW_SQ) index type via vendored faiss GpuIndexHNSW: IndexNode wiring, index-type registration, faiss_gpu_hnsw CUDA target, and recall/cosine/topk/serialize tests. Clean feature-only commit on current main; dead cooperative-distance code excluded. Signed-off-by: Devin AI <devin-ai-integration[bot]@users.noreply.github.com>
Mirror the 6si/faiss GPU HNSW review fixes into the vendored copy: - default overflow_factor to 2 in internal GpuHnswSearchParams - propagate config_.device into the scratch pool (no hardcoded device 0) - make setSearchParams sticky (no consume-once race) - query device shared-memory opt-in limit and opt the kernel into >48KiB dynamic shared memory instead of assuming a 48KiB cap - synchronize before the stack-local entry-point buffer is destroyed Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
- keep Count()/Dim() working after the CPU copy is freed by capturing n_rows/dim before reset and overriding Count()/Dim() (previously returned -1, making the loaded segment look empty) - make HasRawData()/GetVectorByIds() and StaticHasRawData() consistent: GPU_HNSW frees the CPU copy so it exposes no CPU-reconstructable raw data (return false / not_implemented once uploaded) - set gsp.overflow_factor=2 in the search path so the overflow candidate queue stays enabled (recall) - publish gpu_index_ through an atomic gpu_ready_ flag (release/acquire) to remove the unsynchronized double-checked read - harden the filtered-search guard to reject on bitset.data()!=nullptr instead of count() (which defaults to 0) - correct the class comment: GPU_HNSW is registered for F32 and INT8 only - add Count()/Dim() regression assertions to the GPU HNSW search test 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>
Mirror of 6si/faiss#1: the overflow candidate queue was never wired in (overflow_insert had no call site; d_overflow_count only zeroed+read), so it only wasted scratch VRAM + a memset per search. Remove the overflow machinery from the vendored faiss GPU kernel and drop the explicit gsp.overflow_factor=2 in faiss_hnsw.cc. Recall is tuned via ef. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
GPU_HNSW_SQ was exposed by Milvus (constant, param spec, tests, GPU routing) but had no Knowhere backend, so creating it hit no impl. Add a thin GpuHnswSQIndexNode subclass of GpuHnswIndexNode that only overrides Type(); the GPU search path is identical (a CPU-built HNSW/HNSW_SQ index uploaded to GpuIndexHNSW). Register it for VECTOR_FLOAT and VECTOR_INT8, add the INDEX_GPU_HNSW_SQ constant + index-table entries, and accept the GPU_HNSW_SQ binset key in Deserialize's alias path. 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>
Sync of 6si/faiss@467593eae (byte-identical vendored copy): - searchImpl_ event-based stream sync before reading host queries (zilliztech#15) - cudaSetDevice() before cudaFree in scratch/index destructors (zilliztech#19/zilliztech#20) - clear error when search_width exceeds shared-memory budget (zilliztech#18) 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>
Mirror the faiss native fp16/bf16 device-storage change into the vendored faiss copy (byte-identical to 6si/faiss) and register fp16/bf16 for both GPU_HNSW and GPU_HNSW_SQ: - index_table.h: add VECTOR_FLOAT16 / VECTOR_BFLOAT16 legal-type rows for both GPU index types. - faiss_hnsw.cc: add fp16/bf16 KNOWHERE_REGISTER_GLOBAL (feature::FP16|GPU, feature::BF16|GPU) and RegisterStaticFunc entries for both nodes; update the class comment to describe native 2-byte fp16/bf16 device storage with per-element up-conversion to fp32 in the kernel. - test_gpu_search.cc: add GPU HNSW FP16/BF16 deserialize+search UTs (CPU HNSW_SQ QT_fp16/QT_bf16 build -> GPU load -> recall vs brute force). Compile-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>
… bound) - Deserialize (reload path) now clears gpu_ready_ under gpu_mutex_ before resetting gpu_index_, so a failed re-upload no longer leaves the node advertising gpu_ready_==true with a null gpu_index_ (the lock-free reader in Search() would otherwise dereference null). - Round-trip self-recall test searched train_ds (nb rows) but only checked the first nq of nb results; loop and denominator now use nb. - Document the single-visible-GPU-per-process device model of the shared StandardGpuResources (multi-GPU-per-process device selection is a follow-up). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
The existing SQ8 test builds an unsigned QT_8bit HNSW_SQ, which routes through the decode-to-fp32 upload branch, so it does not exercise the native signed-int8 device path (upload_int8_dataset / QT_8bit_direct_signed). Add an int8 test that builds a CPU int8 HNSW and loads it as GPU_HNSW, covering all four native device representations (fp32/int8/fp16/bf16). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…eview fixes - M1: GpuHnswIndexNode::Deserialize now returns Status::cuda_runtime_error when the eager GPU upload throws, instead of returning success with no GPU index (which only surfaced the failure on first search). Also log via LOG_KNOWHERE_ERROR_ rather than fprintf(stderr,...) so it lands in Milvus's structured logs. - Mirror the byte-identical vendored faiss changes (H1 scratch-release sync, L1/L2/C1/H5 documentation) to keep 6si/faiss and the vendored copy in lockstep. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…oad accounting Codex P1 review fixes for the GPU HNSW path: #1 Quantized-cosine inverse-norm semantics. The CPU cosine index records inverse L2 norms from the ORIGINAL input vectors (HasInverseL2Norms:: get_inverse_l2_norms()). The GPU upload previously recomputed norms from decoded/quantized codes (and normalized decoded fp32 vectors in place), which diverges from the CPU scores/traversal for lossy SQ8/int8/fp16/bf16. Now upload the CPU index's stored inverse norms and apply them in the kernel; only flat-fp32 cosine (decoded==original) keeps normalizing in place. Vendored faiss copy kept byte-identical with 6si/faiss. #2 Reload/search use-after-free race. gpu_index_ is now an atomic<shared_ptr<GpuIndexHNSW>>; Search() takes a shared-ownership snapshot for the whole search so a concurrent Deserialize() reload that resets/rebuilds the member cannot free an in-use index. Atomic readiness only guarantees visibility, not lifetime. #3 Phase-separated load-resource accounting. Resource gains maxMemoryCost (transient peak during load; defaults to 0 so other indexes keep their existing heuristic). GpuHnswIndexNode reports memoryCost=0 (CPU copy freed after upload) and maxMemoryCost covering the download buffer + deserialized CPU index + decode/graph staging. Tests (tests/ut/test_gpu_search.cc, [gpu_hnsw_p1]): native low-precision device paths (fp16/bf16/sq8 via non-mock fp32 node), SQ8/fp16/bf16 cosine CPU-vs-GPU parity asserting per-query top-1 id + per-result score on same-direction/ different-magnitude vectors, search-vs-concurrent-reload lifetime stress, >4 concurrent searches with varying nq/k/ef vs own ground truth, and load-resource estimate assertions. 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>
std::atomic<std::shared_ptr<T>> requires C++20 / libstdc++-12; the build toolchain is GCC 11 (libstdc++-11), where it fails the trivially-copyable static_assert. Keep the same lifetime guarantee with a plain shared_ptr guarded by the existing gpu_mutex_: Search() copies it into a local snapshot under a short lock (not held during the search), and Deserialize() already mutates it under the same lock. Reload still cannot free an in-use GPU index. 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>
L40S measures GPU HNSW cosine recall ~0.98 (comparable to L2), so the 0.65 floor was far too loose to catch a real regression. Align it with the L2/IP/TopK sections (0.85-ish) at 0.80. 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>
Add a test-only GpuHnswUploadFaultInjection hook (host-safe, in GpuHnswTypes.h) that arms the Nth wrapped CUDA call in the device-upload path to simulate a failure. New [gpu_hnsw_p1] section forces the eager upload to fail both immediately and mid-upload, and asserts: load returns an error, no half-published index (a re-armed search errors cleanly, no crash), VRAM returns to baseline after teardown (no leak), and a fault-free retry loads + searches correctly. Vendored faiss copy mirrors the faiss change byte-identically. 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>
Mirrors faiss byte-identically: node_ids is read-only after construction, so cudaMemcpy reads it directly. Behavior-neutral. 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>
…path Milvus loads sealed vector-index segments through VectorMemIndex::LoadFromFile -> index_.DeserializeFromFile(), not Deserialize(BinarySet). GpuHnswIndexNode only overrode Deserialize(BinarySet), so the file/mmap load path fell through to the base CPU loader and never ran the eager GPU upload: the CPU-built HNSW stayed resident in host RAM (fp32-expanded for int8/fp16/bf16 inputs, ~5.9 GiB/segment for the mpd_v2 collection) while VRAM sat near-idle, OOMKilling querynodes during the load storm before any search could trigger the lazy-upload fallback. Override DeserializeFromFile to load the CPU index via the base and then run the same eager upload + host-copy release. The shared upload/unpublish logic is factored into UploadCpuIndexToGpuLocked()/UnpublishGpuIndexLocked() so both entrypoints behave identically; GpuHnswSQIndexNode inherits the fix. Adds a [gpu_hnsw_p1] section that loads via DeserializeFromFile (enable_mmap true/false) and asserts Size()==0 before any search (i.e. indexes[0] freed), which fails pre-fix and cannot be masked by the first-search lazy upload. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…p eligibility Two coupled fixes for the mpd_v2 host-RAM OOM regression that appeared once the HNSW->GPU_HNSW load override was introduced. 1. StaticEstimateLoadResource previously proxied the resident deserialized index as ~file_size. For a compressed (int8) on-disk index this grossly under-counts the fp32-expanded hnswlib graph materialized in host RAM, so the loader admitted far too many concurrent uploads and OOMed the host cgroup before any GPU upload completed (VRAM stayed at 0). The estimate now sizes the transient peak as file + resident fp32 vectors (rows*dim*4) + graph neighbor lists (rows*M*2*4) + fp32 decode staging (rows*dim*4), reads M from the load config when present, and falls back to a conservative multiple of file_size when row/dim metadata is missing. 2. GPU_HNSW / GPU_HNSW_SQ were not advertised as mmap-capable, so the override HNSW->GPU_HNSW silently flipped enableMmap off and the whole CPU index landed in anonymous host RAM instead of being file-backed as plain HNSW was before the override. Re-add feature::MMAP to the registrations and list the types in legal_support_mmap_knowhere_index; the CPU index is deserialized via the mmap-capable file path and freed after the VRAM upload. Strengthen the [gpu_hnsw_p1] load-resource test to assert the estimate tracks the fp32 resident footprint (not file_size) for a compressed index and does not collapse when num_rows arrives as 0. 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>
The HNSW->GPU_HNSW load deserializes the on-disk faiss index (IndexHNSWSQCosine with QT_8bit_direct_signed for an int8 collection) into a compact host copy, uploads it to VRAM, then frees it. That transient int8 copy is ~1 byte/dim, matching how the CPU HNSW int8 node loads -- so GPU_HNSW does not need to be memory-mapped. Remove feature::MMAP from the GPU_HNSW/GPU_HNSW_SQ registrations and the mmap support table; do not rely on enable_mmap. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
… fp32 The prior StaticEstimateLoadResource sized the transient host peak as an fp32/hnswlib index (rows*dim*4 resident + rows*dim*4 decode staging + graph), yielding ~20 GiB/seg. mpd_v2's on-disk index is the compact FAISS int8-SQ form (IndexHNSWSQCosine): it deserializes to ~file_size in host RAM and the signed int8 codes are uploaded to the device directly, so there is no fp32 reconstruct/decode staging on this path. The CPU HNSW cluster loads the same index resident at ~file_size (~1.1 GiB/seg measured). Because the estimate is charged into the querynode's committed-loading memory (indexLoadingMemoryContribution), the fp32 overestimate made the soft-OOM admission guard trip after ~3 segments (3x20 > threshold) even though real RSS was only ~16 GiB and freed back to baseline after every upload. Sizing the transient at ~2x file_size (serialized read buffer + deserialized compact index) makes admission reflect the true compact footprint. Retained cost stays 0 (data lives in VRAM after the CPU copy is freed). 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>
The test was written for the intermediate 4d3d97d commit which used an fp32-expansion-based estimate. After 7d8a448 changed the implementation to file_size*2, the test assertions were not updated and would fail on any CUDA-enabled build. Update the int8-file and no-rows test cases to assert >= 2*file_size instead of the old fp32-based bounds. Also clarify the StaticEstimateLoadResource comment to note that the no-fp32-staging claim applies to native int8/fp16/bf16 paths only, not unsigned SQ8 (QT_8bit) which decodes via sa_decode(). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: Premal Shah <premal@6sense.com>
- Add searchHostInt8() to GpuIndexHNSW: applies +128 query shift to match upload_int8_dataset encoding, uploads int8 queries to d_queries_i8 - Add layer0_beam_search_kernel_int8: uses __dp4a for 4x int8 MADs/cycle - Dispatch in GpuHnswIndexNode::Search when data_format==int8, bypassing IndexNodeDataMockWrapper<int8> fp32 upcast - dim=384: max accumulator 6.2M << INT32_MAX, no overflow risk
… type Add INT8 path to gpu_hnsw_search() that dispatches to layer0_beam_search_kernel_int8 when sc.d_queries_i8 is populated. Upper-layer search still uses fp32 queries (d_queries) for the greedy entry-point descent through sparse upper layers.
V5-C1: Remove +128 query shift in searchHostInt8. upload_int8_dataset already reverses FAISS's +128 bias (codes[i]-128), so GPU stores original signed user values. Queries arrive as the same signed int8 values — no shift needed. The +128 shift was causing (q+128)*d instead of q*d. V5-C2: Gate DP4A dispatch on idx.use_ip in GpuHnswSearch.cuh. The DP4A kernel only computes dot products; L2 now falls through to the existing fp32 generic kernel. V5-H1: Replace throw on dim%4!=0 with fp32 fallback via searchHost(). V5-L1: Fix misleading comments describing the encoding.
…l merge) Re-vendor faiss/gpu from 6si/faiss (devin/gpu-hnsw-int8-dp4a-unify) so the vendored copy stops diverging from upstream faiss. This pulls in: - the single <DataT,QueryT,USE_DP4A> layer-0 kernel (replaces the separate serial int8 kernel + serial fp32 kernel with one parallel bitonic-merge kernel that also handles native int8 DP4A); - staging padded to next_pow2(search_width*2*M) so non-power-of-two 2*M no longer hard-fails GPU search; - ef-clamp warning and corrected int8/negated-distance header docs. Also: - document the fp32-Flat/unsigned-SQ8 transient-peak under-count in StaticEstimateLoadResource (2x is correct for the compact int8 path production uses; a per-DataType estimator is the proper fix); - add a non-power-of-two 2*M UT covering both the fp32 and int8 DP4A parallel-merge paths (cosine, dim divisible by 4). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Re-vendors faiss gpu-hnsw-faiss@81c9ff25 GpuHnswSearchKernel.cuh. Corrects the parallel_merge_into_result diagonal binary search (staging[mid] vs result[pos-mid-1], not result[pos-mid]) that caused ~4% recall and a UINT32_MAX-sentinel CUDA illegal memory access on int8/COSINE search. Validated on CPU vs std::merge (200k cases, 0 failures). All 7 vendored faiss/gpu files remain byte-identical to faiss. Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
… tests Re-vendor faiss GpuHnswSearchKernel.cuh (faiss gpu-hnsw-faiss @46235ee3): guard ep/parent/result node ids < N before graph/dataset indexing (fixes v82 fp16 P1 OOB in layer0_beam_search_kernel<__half,float,false>). Add independently-taggable per-dtype GPU HNSW search tests (nb=4000, dim=64, M=16, ef=200 -- same config as the failing fp16 P1 case), each covering L2/IP/COSINE so a CUDA context crash in one instantiation does not hide the others: - [gpu_hnsw_dtype_fp32] flat -> <float,float,false> (config control) - [gpu_hnsw_dtype_fp16] HNSW_SQ fp16 -> <__half,float,false> - [gpu_hnsw_dtype_bf16] HNSW_SQ bf16 -> <__nv_bfloat16,float,false> - [gpu_hnsw_dtype_sq8] HNSW_SQ sq8 -> decoded <float,float,false> - [gpu_hnsw_dtype_int8] native int8 (IP/COSINE take DP4A, dim%4==0) 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>
Re-vendor faiss GpuHnswSearchKernel.cuh (faiss gpu-hnsw-faiss @81e4cc69): split the bitonic compare-exchange into read -> __syncthreads -> write -> __syncthreads phases, removing the cross-warp read/write hazard that produced the v82 fp16 P1 OOB. Vendored file byte-identical to faiss. 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>
…OSINE re-negation; re-vendor faiss Source fixes (src/index/hnsw/faiss_hnsw.cc): - GPU_HNSW_SQ deserialize: alias the incoming payload under this node's own Type() (INDEX_GPU_HNSW_SQ) so the base loader's GetByName(Type()) finds it; the previous hardcoded INDEX_GPU_HNSW alias made every GPU_HNSW_SQ segment unqueryable while leaving base GPU_HNSW behavior intact. - Lazy GPU upload now reuses UploadCpuIndexToGpuLocked() (made const) instead of deriving the metric from the query config, so IP/COSINE/L2 is detected from the FAISS index even when the query omits metric_type. - Remove the IP/COSINE sign-flip loops after searchHost()/searchHostInt8(): faiss now returns true similarity on copy-out, so re-negating would invert it. - Correct the stale native-int8 comments (no +128 shift; signed int8 uploaded verbatim, DP4A kernel). Re-vendored faiss GPU tree (thirdparty/faiss, byte-identical to faiss 48ab2cb7): GpuIndexHNSW.h, impl/GpuHnswSearch.cuh, impl/GpuHnswSearchKernel.cuh, impl/GpuHnswTypes.cu. Tests (tests/ut/test_gpu_search.cc): - [gpu_hnsw_sq_alias] load a CPU HNSW_SQ index into the GPU_HNSW_SQ node. - [gpu_hnsw_ip_sign] IP/COSINE distances are true similarity (sign+magnitude vs recomputed neighbor similarity; descending order). - [gpu_hnsw_k_valid] all k slots are real neighbors (no sentinel) with ef>=k. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…lute floor The [gpu_hnsw_compare] 'Varying ef comparison' section asserted gpu_recall >= 0.5f for every ef, but at ef=16 (~= k) HNSW recall is legitimately low on both engines (CPU HNSW itself scores ~0.41, GPU ~0.47). The absolute floor was meaningless in that regime and the section already computes cpu_recall but never used it. Assert GPU tracks the CPU oracle within a small tolerance instead, which is the correct invariant for a GPU-vs-CPU comparison test and catches real regressions at every ef. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…nfig I1: GpuHnswIndexNode::Search() detected COSINE from hnsw_cfg.metric_type to decide fp32 query normalization and int8 1/||q|| distance scaling. A search config that omits metric_type defaults to L2, so a cosine index would skip normalization and silently return wrong scores. Derive the metric from the built GPU index via the new GpuIndexHNSW::isCosine() accessor (captured at upload time from the FAISS index type), matching the index-side detection already used on the upload path. Re-vendor the faiss GPU tree (byte-identical to 6si/faiss@24be1b66): - mutex-guarded, monotonic cudaFuncSetAttribute smem configuration (TOCTOU fix) - entry_point bounds validation in extract_hnsw_layers - GpuIndexHNSW isCosine()/useInnerProduct() accessors Add a regression test that searches a COSINE index with metric_type omitted from the config and asserts identical results to the explicit-COSINE search. 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>
…lags) Byte-identical re-vendor of 6si/faiss@a43d3afe: reset() now clears is_cosine_/use_ip_ so a reused GpuIndexHNSW can't inherit a stale cosine/IP interpretation. Verified with cmp against the faiss source. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
The faiss_gpu_hnsw OBJECT library is the only place the faiss GPU sources are compiled in the knowhere build: FAISS_SRCS does not glob faiss/gpu/*, and the vendored thirdparty/faiss/faiss/gpu/CMakeLists.txt (faiss_gpu_objs) is never add_subdirectory()'d. Comment-only; no compiled objects 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>
Re-vendor the faiss GPU tree with CPU-parity filtered search and wire the delete/TTL/partition BitsetView through GpuHnswIndexNode::Search() instead of rejecting a non-empty bitset. - faiss thirdparty: byte-identical copy of the filtered-search kernel, GPU brute-force header (GpuHnswBruteForce.cuh, added to gpu/CMakeLists.txt), host launcher and scratch types. - Search(): remove the blanket filtered-search rejection; pass bitset ptr/nbits/filtered-count and disable_fallback_brute_force into faiss for both native int8 and fp32 paths. Still reject id-remapped / id-offset bitsets (has_out_ids()/id_offset()!=0) since the kernel tests by raw node id (== row offset == bit index), and guard a bitset shorter than ntotal. Snapshot/reload concurrency, CPU-index release after upload, and cosine/int8 post-scaling preserved. - bitsetview.h: add id_offset() accessor for the shape guard. - tests: gpu_hnsw_id_mapping (blocking G-ID proof: node id == row offset == bitset index, via unfiltered self-query and filtered self-query) and gpu_hnsw_filtered_recall (no returned id filtered; recall vs brute force across delete ratios 0.1-0.99). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…edge cases; vendor ensure_filter device fix Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
feat(gpu_hnsw): GPU-side filtered search (deletes/TTL/partition bitset)
…stance, H1) Re-vendors faiss gpu-hnsw-faiss tip f22116e8 (merge of faiss PR #8): the layer-0 neighbor expansion is now warp-per-candidate cooperative distance (32 lanes stride one vector, warp-reduce) instead of thread-per-candidate, coalescing the dominant dataset loads. Measured on the L40S build box (counter deltas transfer to Blackwell): 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.0025), compute-sanitizer clean (memcheck/racecheck/initcheck). No public API or semantic change. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
feat(gpu_hnsw): re-vendor faiss GPU tree (warp-cooperative layer-0 distance, H1)
…OOM fix) Re-vendors the faiss GPU HNSW fix that bounds the grow-only visited bitmap by processing queries in nq-chunks (byte-identical to 6si/faiss devin/gpu-hnsw-oom-chunk). The bitmap was nq * ceil(N/32) * 4 bytes per scratch slot and never released, so high search concurrency on large segments exhausted device VRAM. Chunking caps it regardless of batch size / concurrency; small segments still run a single pass. Adds tests/ut/test_gpu_search.cc [gpu_hnsw_chunk_parity]: runs each search twice (default cap => single pass, vs GPU_HNSW_BITMAP_BYTES forced tiny => many chunks) and asserts identical ids/distances across fp32 & int8-DP4A, unfiltered, mid-ratio filtered, and up-front-BF (99%) paths for L2/IP/COSINE. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
feat(gpu_hnsw): re-vendor faiss GPU tree (nq-chunked visited bitmap, OOM fix)
…uIndexHNSW Re-vendor the faiss gpu-hnsw-faiss-native tree (vanilla-faiss-structured GpuIndexHNSW built on faiss::IndexHNSW, cppcontrib-free) and rewire GpuHnswIndexNode to be a pure consumer of it: - GpuHnswBuild.cuh (cppcontrib-coupled) -> GpuHnswBuildCommon.cuh + GpuHnswBuildVanilla.cuh; GpuCloner/GpuAutoTune/SWIG updated (knowhere cuVS-disable patches preserved). - UploadCpuIndexToGpuLocked converts the deserialized cppcontrib::knowhere::IndexHNSW to a vanilla faiss::IndexHNSW (ToVanillaHnsw: field-copy graph, borrow Flat/SQ storage for L2/IP, reconstruct+normalize+re-encode for cosine=normalize+IP) and uploads via copyFrom() instead of copyFromWithMetric(). - Track cosine on the node (gpu_is_cosine_) since the GPU index only knows L2 vs IP; query-side normalization unchanged. Quantized-cosine numerics differ from the legacy inverse-norm path (gated by recall tests; may require re-ingest). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
QT_8bit_direct_signed is a fixed code=x+128 map with no trained range, so L2-normalized cosine vectors (components ~1/sqrt(d)) collapse onto a handful of the 256 levels and int8-cosine recall is destroyed. ToVanillaHnsw now re-encodes int8 cosine as fp16 (accurate at any dim, validated fp16 GPU path); the raw int8 query is widened to fp32 and routed through searchHost, not the DP4A path. int8 L2/IP are unchanged (direct_signed + DP4A). The quantized-cosine test now gates GPU recall against an exact brute-force cosine oracle (representation-agnostic) instead of asserting byte/score parity with the legacy inverse-norm cppcontrib path, which the native normalize+IP path deliberately abandons. 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>
…f16 vs oracle Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Author
🤖 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:
|
devin-ai-integration
Bot
changed the title
![devin-ai-integration[bot]](https://avatars.githubusercontent.com/in/811515?s=60&v=4){kind=small}
Mirror the two Devin Review fixes from faiss PR #10 into thirdparty/faiss so the built libknowhere matches: - upload the filter bitset in GpuIndexHNSW::searchImpl_ before gpu_hnsw_search (the standard search() path selected the filtered kernel on a sticky-param bitset but never allocated/uploaded the device buffer -> illegal access). - guard extract_hnsw_layers against reading nb_neighbors(1) for a degenerate layer-0-only index (max_level==0), which would be an out-of-bounds read of cum_nneighbor_per_level[2]. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Author
devin-ai-integration Bot
added a commit
that referenced
this pull request
Jul 22, 2026
Mirror faiss #11 review fix into the vendored faiss GPU tree: select the native int8 DP4A layer-0 path on a per-search i8_queries_staged flag (set in GpuHnswSearchScratch::ensure from use_i8_queries) instead of sc.d_queries_i8 != nullptr. Pooled scratch slots keep d_queries_i8 allocated after an int8 search, so a later fp32-query search on the same slot could take the DP4A path against stale int8 query data; gating on the staged flag makes the fp32 fallback correct for reused slots. Also wraps vendored comment/source lines to 80 chars (no behavior change). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Adds knowhere GPU HNSW (
GPU_HNSW,GPU_HNSW_SQalias) backed by the faiss-nativefaiss::gpu::GpuIndexHNSW(see faiss PR), across all dtypes/metrics, with GPU-side filtering and accurate int8 cosine. This branch is the whole feature intomain. On-disk format is unchanged — existing serialized segments load; the cppcontrib CPU index is converted to a vanillafaiss::IndexHNSWin memory before GPU upload.Index node (
src/index/hnsw/faiss_hnsw.cc)ToVanillaHnsw(...)builds a vanillafaiss::IndexHNSWfrom the loaded index and uploads viaGpuIndexHNSW::copyFrom(...)(drops fork-onlycopyFromWithMetric/isCosine()).GetGpuConstructionMutex(), mutex-guarded reload-safeshared_ptrsnapshots, bitset filtering, memory admission (no CPU fallback), GPU-upload-failure → failDeserialize.Accurate int8 cosine (substantive fix)
QT_8bit_direct_signed(fixedcode=x+128) collapses L2-normalized vectors onto a few of 256 levels. Fix is entirely host-side at load (datanode RAM, not querynode VRAM):Query path: DP4A dispatch guarded
if (int8 && !is_cosine). Int8 cosine queries widen int8→fp32, normalize in place, runsearchHost()against the fp16 DB (symmetric); int8 L2/IP keepsearchHostInt8()/DP4A. Cost: 2 B/dim (fp16) for int8-cosine segments only; a 1×-VRAM trained-QT_8biton-the-fly-decode path is a deferred follow-up.GPU-side filtering + resource accounting
Bitset filters (deletes/TTL/partition/visibility) applied on-device with CPU-parity recall; load-estimate sized to the compact int8 transient (~2× file) not fp32; GPU upload runs on the
DeserializeFromFile(mmap) path; peak host-RAM reserved under tiered eviction.Tests (
tests/ut/test_gpu_search.cc+ dtype/filtered suites)Full dtype × metric matrix; filtered recall vs CPU-HNSW baseline; ef-sweep GPU-vs-CPU parity. Quantized-cosine gate rewritten to a brute-force cosine oracle: fp16/bf16 keep absolute bars (0.85/0.80) + GPU-tracks-CPU; sq8 only requires
gpu_recall >= cpu_recall - 0.02(a single trained range can't preserve direction across the adversarial 1000× magnitude spread).Vendored faiss
thirdparty/faisstracks the faiss PR's GPU HNSW tree.Validation (CUDA, GPU CI helper)
All 16 GPU HNSW tests: 202,360 assertions, 0 failures. Recall: int8-COSINE 0.9975, sq8-COSINE 0.9855, fp16-COSINE 0.998, bf16-COSINE 0.983.
Link to Devin session: https://6sense.devinenterprise.com/sessions/55dab0bd33c346df99b5818b30059342
Requested by: @premal