feat(gpu): GPU HNSW for faiss — GpuIndexHNSW (int8/DP4A, fp16/bf16, filtered, cloner-bridged) by devin-ai-integration[bot] · Pull Request #10 · 6si/faiss

devin-ai-integration · 2026-07-21T13:32:51Z

Summary

Adds GPU HNSW to faiss as a first-class GPU index. GpuIndexHNSW is built on vanilla faiss::IndexHNSW (Flat/SQ storage + faiss::HNSW graph) and produced by the standard cloner, exactly like GpuIndexFlat/GpuIndexIVF*:

   CPU                    GPU
   IndexHNSW  --cpu_to_gpu-->  GpuIndexHNSW   (search-only; copyTo throws)

This branch is the whole feature (single squashed history into main); it evolved from an OCQ-beam prototype to the current native design. Highlights below.

Index & cloner

GpuIndexHNSW (.h/.cu): copyFrom(const faiss::IndexHNSW*) maps faiss::HNSW (CSR neighbor arrays, entry_point, levels, cum_nneighbor_per_level) to a flat device graph and uploads storage. GPU→CPU (copyTo/index_gpu_to_cpu) intentionally throws (search-only). Cosine = normalize + METRIC_INNER_PRODUCT (no COSINE metric / inverse-norm buffer), matching faiss GPU convention.
GpuCloner.cpp: CPU→GPU routes faiss::IndexHNSW — excluding IndexHNSWCagra (stays on cuVS/CAGRA) — to GpuIndexHNSW.
SWIG (DOWNCAST_GPU(GpuIndexHNSW)) + AutoTune (efSearch sweep).

Search kernel (perf + correctness)

Unified layer-0 kernel with native int8 DP4A dot-product for QT_8bit_direct_signed and warp-cooperative, coalesced dataset loads.
Native FP16/BF16 device storage (QT_fp16/QT_bf16) — half-width loads, no fp32 expansion.
Parallel bitonic-sort + merge-path result merge (replaced serial merge); fixed a merge-path off-by-one and a bitonic_sort_staging shared-memory race.
CPU-parity filtered search: bitset filter (deletes / TTL / partition / visibility) applied on-device; per-GPU device binding before cudaMalloc.
VRAM safety: nq-chunked layer-0 search bounds the visited-bitmap allocation; OOB guards on entry-point/parent/result node ids; thread-safe cudaFuncSetAttribute smem config.

int8 accuracy (paired with the knowhere consumer)

INT8 L2/IP use QT_8bit_direct_signed + DP4A. There is deliberately no SQ_Int8_Cosine gate: direct_signed is a fixed code=x+128 map, so L2-normalized vectors (components ~1/sqrt(d)) collapse onto a few of 256 levels. The knowhere consumer re-encodes int8-cosine as fp16 before upload; SQ_Fp16_Cosine is the representative gate.

Tests (`TestGpuIndexHNSW.cpp`)

Flat L2/IP/cosine, SQ int8 L2 (int8-range data), SQ fp16 L2, SQ fp16/bf16 cosine; cloner dynamic type, valid IDs, distance sign/order, recall, unsupported copyTo, unsupported metric.

Validation (CUDA arch 89, GPU CI helper)

TestGpuIndexHNSW 9/9 PASS. Consumed downstream by knowhere (int8-COSINE recall 0.9975; 202,360 assertions, 0 failures). Image milvus-gpu:gpu-hnsw-faiss-native-v1 (sha256:9deb6299…) built from this stack.

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

GPU HNSW index for faiss (GpuIndexHNSW): build + OCQ beam-search kernel, int8/fp16/fp32 support. Clean feature-only commit on current main; dead cooperative-distance code (select_threads_per_dist/coop_l2_distance/ coop_ip_distance, zero callers) excluded. Signed-off-by: Devin AI <devin-ai-integration[bot]@users.noreply.github.com>

- default overflow_factor to 2 in internal GpuHnswSearchParams so the default search path keeps the overflow candidate queue enabled - propagate config_.device into the scratch pool instead of hardcoding device 0 (fixes cross-device streams on non-default devices) - make setSearchParams sticky (no consume-once race under concurrency) - align copyFrom header signature with the knowhere::IndexHNSW impl and declare copyFromWithMetric - query the 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>

The overflow candidate queue was never wired in: overflow_insert() had no call site and d_overflow_count was only ever zeroed and read, so the num_parents==0 fallback always no-oped. It only allocated nq*overflow_factor*ef scratch VRAM + a per-search memset for a queue that was never populated. Remove overflow_factor from the public/internal search params, the d_overflow_* scratch buffers, overflow_insert(), and the dead fallback. The kernel is a plain parallel beam search; raise ef for more recall. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

…em clamp) - searchImpl_: make the scratch-pool stream wait on the resources' default stream (event) before reading queries copied there by GpuIndex::search, fixing a host-query race (finding facebookresearch#15). - GpuHnswSearchScratch/GpuHnswDeviceIndex: record the owning device and cudaSetDevice() before cudaFree in their destructors (multi-GPU, facebookresearch#19/facebookresearch#20). - GpuHnswSearch: throw a clear error when search_width is too large to fit even ef=1 in shared memory instead of clamping ef negative (facebookresearch#18). Findings facebookresearch#17 (level off-by-one) and facebookresearch#21/facebookresearch#22 (IP/cosine sign) are false positives: this fork uses the standard FAISS levels[] convention and Knowhere un-negates IP/COSINE and normalizes cosine queries around searchHost. 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 native fp16/bf16 support to GpuIndexHNSW, mirroring the existing int8 native path. fp16/bf16 HNSW_SQ codes (QT_fp16/QT_bf16) are stored raw as IEEE half/bfloat16 (2 bytes/element), bit-compatible with CUDA half / __nv_bfloat16, and uploaded verbatim to the device (no up-conversion to fp32 at upload, preserving the memory win). The search kernel dispatches on a new GpuHnswDatasetType enum (FP32/INT8/FP16/BF16) and the templated distance helpers select the matching load_elem overload, up-converting each element to fp32 for accumulation. Cosine inverse-L2-norms are computed from the fp32-decoded values (mirroring int8). Graph-walk algorithm unchanged. Compile-verified only; GPU-runtime validation pending capacity. 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>

…strings Address external review findings on the GPU HNSW code: - H1: add cudaStreamSynchronize safety net in GpuHnswScratchPool::release so a slot's stream is drained before the slot is reused; guards against a caller that skips its own sync (latent use-after-free via ensure() free+realloc). - L2: GpuIndexHNSW docstring updated to list fp16/bf16 (native 2-byte device storage) alongside float32/int8, and copyFrom SQ type list. - C1: document that searchHost is the preferred entry point and searchImpl_ round-trips labels D2H->H2D; a GPU-side label-conversion kernel is a follow-up. Knowhere uses searchHost. - H5: document the HnswT levels[] 1-based-count convention required by extract_hnsw_layers. - L1: document num_layers as informational (search uses num_upper_layers_built). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

The CPU cosine index records inverse L2 norms from the ORIGINAL input vectors (HasInverseL2Norms::get_inverse_l2_norms()). Recomputing norms from decoded/ quantized codes (or normalizing decoded fp32 vectors in place) diverges from the CPU index's scores and HNSW traversal for lossy SQ8/int8/fp16/bf16 cosine. Thread the stored inverse norms from from_faiss_hnsw_sq into the int8/halfwidth/fp32 upload paths and apply them in the kernel; only flat-fp32 cosine (decoded == original) keeps normalizing in place. Byte-identical with the knowhere vendored copy. 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 GpuHnswUploadFaultInjection (host-safe, in GpuHnswTypes.h): a countdown that makes the Nth GPU_HNSW_BUILD_CUDA_CHECK in the device-upload path report a simulated failure. Never armed in production (countdown 0 -> one relaxed atomic load per wrapped call, no effect); lets Knowhere unit tests exercise the upload error/cleanup/retry path without a real OOM. 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>

node_ids is read-only after construction, so copy directly from it in the cudaMemcpy instead of duplicating into h_node_ids. 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>

Phase 1 optimization: all 128 threads now participate in both the staging sort and result merge instead of thread 0 doing all the work while the other 127 idle. Changes: - GpuHnswSearchKernel.cuh: add bitonic_sort_staging (128-element, 128-thread, in-place) and parallel_merge_into_result (binary-search per output slot, O(log N) divergence) as replacements for the serial sorted-insert loop; add merged_ids/merged_dists/merged_expanded scratch arrays to shared memory; update calc_layer0_smem_size - GpuHnswSearch.cuh: update smem_overhead formula (24 bytes/ef slot instead of 12, adds 3 merge scratch arrays); add power-of-2 check on search_width × max_degree0 required by bitonic sort Three correctness bugs fixed during development: 1. Bitonic sort race: each thread now writes only to slot i (not partner), eliminating inter-warp write-write races (778K instances detected by compute-sanitizer --racecheck) 2. Negative ri index: binary search exit can yield si = min(pos,sc)+1 giving ri = pos-si = -1; added ri_valid guard and ri_safe = max(0,ri) clamp on all array reads 3. Speculative -O2 loads: nvcc speculatively evaluates array loads even when the branch predicate is false; fixed by clamping ri_mid = max(0, pos-mid) in the binary search body and adding an explicit sentinel override when pos < mid Measured results (milvus-gpu-hnsw cluster, mpd_v2 570M INT8, 8 querynodes, v76 serial baseline vs v78 this commit): Single client: batch=256: 4,638 vec/s vs 4,022 (+15%), p50 52ms vs 63ms (-17%) batch=128: 3,783 vec/s vs 3,685 (+3%), p50 33ms vs 32ms (flat) 8 parallel clients (production-representative): batch=64: 11,181 vs 9,919 vec/s (+13%) batch=128: 14,720 vs 9,986 vec/s (+47%) batch=256: 17,432 vs 7,355 vec/s (+137%) ncu barrier stall: 69.89% → 39.67% (-30pp) ncu long-scoreboard stall: 4.77% → 21.01% (+16pp, now new bottleneck) Remaining stall is global memory latency from distance compute (__ldg on INT8 dataset), not the merge — next optimization target. Recall unchanged: R@1=0.899, R@10=0.937 across all configurations. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

… merge Forward-port Knowhere's native int8 DP4A search into faiss's parallel bitonic-merge kernel so the two GPU HNSW trees stop diverging. The two layer-0 kernels (serial int8 + parallel fp32) collapse into one kernel templated on <DataT, QueryT, USE_DP4A>: - generic path: QueryT=float, USE_DP4A=false (fp32/fp16/bf16/int8 via load_elem, L2 or negated-IP); - native int8 path: DataT=QueryT=int8_t, USE_DP4A=true, thread_ip_distance_dp4a. Both share the same shared-memory layout and parallel_merge_into_result. Adds searchHostInt8() (uploads signed int8 queries verbatim, no shift; fp32 fallback when dim%4!=0) plus the d_queries_i8 scratch field/ensure() arg, and INT8 DP4A dispatch (int8 + inner-product/cosine + dim%4==0). Review fixes folded in: - pad layer-0 staging to next_pow2(search_width*max_degree0) and grow the block to match, instead of hard-failing when 2*M is not a power of two; throw a clear error only if the padded capacity exceeds the 1024 block limit; - warn (stderr) when ef is clamped to fit shared memory instead of silently truncating; - document the negated IP/cosine distance convention and the (corrected, no +128 shift) int8 query encoding in GpuIndexHNSW.h. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

…esult The diagonal binary search compared staging[mid] against result[pos-mid] instead of result[pos-mid-1], mis-partitioning every merge (~4% recall on mpd_v2) and, with padded staging, promoting a UINT32_MAX sentinel into the result list that was then used as a graph/dataset index -> CUDA illegal memory access (the v81 crash at GpuIndexHNSW.cu:332). Fix uses ri_mid = pos-mid-1 with a -inf boundary when ri_mid<0. Validated on CPU against std::merge over 200k random (rc,sc,ef) cases: 0 failures (was 187330/200000). Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

…raph indexing Bounds-check every node id before it indexes d_layer0_graph or d_dataset in layer0_beam_search_kernel: - entry point (ep): skip seeding + neighbor scan if ep >= N (ep_valid gate) - parent id: skip expansion if parent >= N (mirrors the existing nbr guard) - top-k output: emit UINT64_MAX sentinel if result_ids[i] >= N Prevents the v82 fp16 P1 OOB (illegal global read in layer0_beam_search_kernel<__half,float,false>, addr ~= bad_id*max_degree0*4) from a stray/garbage node id reaching graph/dataset indexing. On the healthy path (valid ep, correct merge) none of these fire, so recall 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>

…pare-exchange) Split the bitonic compare-exchange into two barrier-separated phases. Each thread owns writes to slot i but READS slot partner=i^j; for partners in a different warp (j >= warpSize) that read raced with partner's write to its own slot in the same phase (racecheck: ~1.6M hazards in <__half,float,false>, ~664K in <float,float,false>). Now: read pair + decide swap -> __syncthreads -> conditional write of own slot -> __syncthreads. This is the actual root cause of the v82 fp16 P1 OOB (the stale read produced a garbage node id); fp32 hit the same race but only OOB'd at large N. Sort result is unchanged (canonical two-phase bitonic); loop bounds stay uniform so every thread reaches every barrier. 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>

…amp ef>=k - Negate IP/COSINE score on kernel copy-out so the public API returns true similarity (larger == closer), matching faiss METRIC_INNER_PRODUCT; make the padding sentinels metric-aware (-FLT_MAX for IP/COSINE, FLT_MAX for L2). - Validate ef>0 and search_width>0 before computing the default iteration count (prevents divide-by-zero). - Auto-clamp ef=max(ef,k) so k results are always tracked; retain the shared-memory-fit guard that throws when k itself cannot fit. - Cache cudaFuncSetAttribute per high-water dynamic shared-memory size. - Set the owning device before destroying a scratch stream (multi-GPU). - Remove unused <atomic> from GpuIndexHNSW.h; correct the public distance documentation; wrap the long merge-path line. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

…pose metric - Replace the racy std::atomic load-check-store around cudaFuncSetAttribute with a mutex-guarded, monotonic high-water update. Two concurrent searches with different smem sizes could interleave so a smaller-ef call ran after a larger one and downgraded the kernel's global max-dynamic-smem attribute while a recorded high-water mark said otherwise, making a later intermediate launch fail. Under the lock the attribute only ever grows and is always >= the current launch's requirement. - Validate hnsw.entry_point in [0, n_rows) in extract_hnsw_layers before it is used to index d_dataset on device, turning a malformed/sentinel entry_point into a host-side throw instead of an OOB device read. - Add GpuIndexHNSW::isCosine()/useInnerProduct() accessors, set at copyFrom()/copyFromWithMetric() time, so callers can post-process results from the index-side metric instead of a per-search config that may omit metric_type (and default to L2). 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>

reset() left is_cosine_/use_ip_ at their prior copyFrom() values. Safe today (searches guard on is_trained), but a subsequent copyFrom() of an L2 index would inherit stale cosine/IP interpretation. Clear them so reset() fully returns the index to an untrained, metric-neutral state. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

…set) Add GPU-side bitset filtering to layer-0 HNSW search so GPU_HNSW serves segments with deletes/TTL/partition visibility, matching CPU HNSW rather than rejecting a non-empty bitset. Kernel: is_bitset_filtered() device test (byte id/8, bit id%8, LSB-first, matches Knowhere BitsetView); filtered rows stay graph waypoints but are never emitted. Two-tier beam (valid ef + invalid frontier ef_inv) with serialized post-sort alpha gate (accumulated_alpha per query in meta[4], init 1.0f; kAlpha=filter_ratio*0.7) and cross-beam parent selection (globally-closest unexpanded across both beams). HAS_FILTER template param; unfiltered path codegen-identical. Brute force (GpuHnswBruteForce.cuh): GPU-resident up-front BF (>=0.93 filtered or k>=0.5*live) and per-query fallback via device needs_bf[] worklist (no host round trip). Host: per-slot d_bitset/d_needs_bf scratch (grow-only ensure_filter), bitset upload before launch, needs_bf_count zeroed via cudaMemsetAsync, smem re-budget (36 B/ef filtered), disable_fallback_brute_force threading. Threaded through all 5 dtypes. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

…GPU) Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

feat(gpu_hnsw): CPU-parity filtered search (deletes/TTL/partition bitset)

The layer-0 expansion computed each candidate's full-dim distance on a single thread, so a warp's 32 lanes read 32 different vectors (rows dim apart) -> ncu measured 12.1 sectors/request (~8% coalescing), 49% of warp stalls on global load latency, and DRAM at only 37.6% of peak. Assign one warp per candidate edge: lane 0 resolves the neighbor id and claims the visited-bitmap first-visit, broadcasts go/no-go + id, then all 32 lanes stride the vector and warp-reduce the partial sums (DP4A int8, generic IP, and L2). Loads coalesce toward 1 sector/request on the dominant traffic source. The predicate is warp-uniform so the warp never diverges across the __shfl reductions. Host rounds block_size up to a whole number of warps. Distance values, staging contents, and result ordering are unchanged (only the parallelization of the identical computation changes); correctness/recall to be confirmed by the existing recall + sanitizer suite. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

perf(gpu_hnsw): warp-cooperative layer-0 distance (coalesced dataset loads)

The layer-0 search scratch is grow-only per pool slot: the visited bitmap is nq * ceil(N/32) * 4 bytes and, once grown, is never released. Under high search concurrency (one bitmap per concurrent pool slot) on large segments this exhausted device memory -- 16 concurrent batch=512 searches on a ~538M-row segment grew the pool to ~97 of ~98 GB and OOM'd every subsequent allocation. Process queries in nq-chunks so the bitmap is bounded regardless of batch size / concurrency. gpu_hnsw_bitmap_chunk() derives the per-launch query count from a VRAM cap (default 256 MiB, env-tunable via GPU_HNSW_BITMAP_MB / GPU_HNSW_BITMAP_BYTES); GpuHnswSearchScratch::ensure() sizes the bitmap for one chunk and the search launches per chunk with caller-offset base pointers. Queries are independent in HNSW search and every per-query buffer is indexed by the chunk-local block index, so chunking is result-invariant (no kernel changes). Small segments yield chunk == nq -> a single pass, identical to prior behavior. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

fix(gpu_hnsw): bound visited-bitmap VRAM via nq-chunked layer-0 search

Port GpuIndexHNSW off Knowhere's cppcontrib layer onto vanilla faiss::IndexHNSW / faiss::HNSW so it fits the standard faiss CPU<->GPU index structure: - copyFrom(const faiss::IndexHNSW*) and a GpuIndexHNSW(provider, const IndexHNSW*, config) convenience ctor, mirroring GpuIndexFlat. - copyTo(faiss::IndexHNSW*) present but throws: the index is search-only (uploads a CPU-built graph, no reconstructable CPU copy). - Register both directions in GpuCloner: index_cpu_to_gpu builds a GpuIndexHNSW from IndexHNSWFlat/IndexHNSWSQ (IndexHNSWCagra excluded, still handled by the cuVS GpuIndexCagra branch); index_gpu_to_cpu throws an explicit unsupported error. - Metric contract: METRIC_L2 and METRIC_INNER_PRODUCT only. Cosine is the standard faiss idiom (normalize + IP), so the Knowhere HasInverseL2Norms / inverse-norm coupling is dropped (d_inv_norms stays null; the kernel already guards on it). - Split the graph/vector upload helpers into GpuHnswBuildCommon.cuh (no cppcontrib dependency; extract_hnsw_layers already matches the faiss::HNSW accessor interface) and add GpuHnswBuildVanilla.cuh with from_index_hnsw_flat / from_index_hnsw_sq. Remove the Knowhere-coupled GpuHnswBuild.cuh from this branch. - Add TestGpuIndexHNSW.cpp: CPU-parity recall gates for Flat L2/IP/cosine and SQ int8/fp16, valid-id and distance-sign checks, and the copyTo/unsupported-metric error paths. Does not touch the production Knowhere-coupled path on gpu-hnsw-faiss. Not built locally (no CUDA toolkit); GPU compile + test run to follow on CUDA hardware. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

- Expose GpuIndexHNSW (and GpuIndexHNSWConfig / SearchParametersGpuHNSW) to Python via swigfaiss: %include the header, add a DOWNCAST_GPU entry so index_cpu_to_gpu() returns a typed GpuIndexHNSW, and %ignore the host-pointer / bitset entry points that take the SWIG-unparsed GpuHnswSearchParams (Python uses the standard search()). - GpuParameterSpace: add an 'efSearch' sweep range for GpuIndexHNSW and route set_index_parameter('efSearch', v) to setSearchParams(), mirroring the CPU IndexHNSW autotune parameter. index_factory needs no change: a factory-built IndexHNSWFlat/SQ clones to GpuIndexHNSW through the GpuCloner registration. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

Cosine over scalar-quantized storage = normalize + IP with the SQ codec. fp16/bf16 represent normalized components (magnitude ~1/sqrt(d)) accurately. int8 cosine has no gate here on purpose: QT_8bit_direct_signed is a fixed code=x+128 map, so normalized vectors collapse onto a few of the 256 levels; the Knowhere consumer re-encodes int8 cosine as fp16, which is SQ_Fp16_Cosine. Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

devin-ai-integration · 2026-07-21T13:32:55Z

🤖 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

devin-ai-integration

Devin Review found 1 potential issue.

View 1 additional finding in Devin Review.

devin-ai-integration · 2026-07-21T22:40:37Z

+        GPU_HNSW_BUILD_CUDA_CHECK(cudaMalloc(
+                &ul.d_neighbors,
+                ul.num_nodes * maxM * sizeof(uint32_t)));
+        GPU_HNSW_BUILD_CUDA_CHECK(cudaMemcpy(
+                ul.d_neighbors,
+                h_neighbors.data(),
+                ul.num_nodes * maxM * sizeof(uint32_t),
+                cudaMemcpyHostToDevice));


📝 Info: Upper-layer neighbor-array byte size computed in 32-bit before widening

In extract_hnsw_layers, allocations such as cudaMalloc(&ul.d_neighbors, ul.num_nodes * maxM * sizeof(uint32_t)) (faiss/gpu/impl/GpuHnswBuildCommon.cuh:160-167) compute ul.num_nodes * maxM in 32-bit (uint32 * int) before promotion to size_t. For upper layers node counts shrink geometrically (~n_rows/M per level), so overflow is very unlikely, but the layer-0 path already guards this by casting n_rows to size_t first (upload_graph_to_gpu, line 204-205). For consistency and safety on extremely large segments the upper-layer products could also cast to size_t before multiplying.

Was this helpful? React with 👍 or 👎 to provide feedback.

devin-ai-integration · 2026-07-21T22:40:58Z

+        if (lane == 0) {
+            int slot = atomicAdd(&meta[1], 1);
+            if (slot < max_staging) {
+                staging_ids[slot] = nbr;
+                staging_dists[slot] = dist;
+            }
+        }


📝 Info: Dropped neighbors when staging overflows are still marked visited

In layer0_beam_search_kernel, when a candidate edge is expanded, bitmap_visit marks the node visited BEFORE the distance is computed, and only lane 0 stores it into staging via atomicAdd(&meta[1], 1) guarded by if (slot < max_staging) (faiss/gpu/impl/GpuHnswSearchKernel.cuh:843-849 and 1000-1006). If a single iteration produces more than max_staging = padded_staging_capacity(search_width, max_degree0) newly-visited neighbors, the overflow candidates are silently dropped yet remain marked visited, so they can never be reconsidered. search_width * max_degree0 newly-visited nodes per iteration is possible when few are already visited, so this can lose true neighbors and slightly reduce recall. Not a crash, and recall gates in the PR still pass, but worth noting for tuning.

Was this helpful? React with 👍 or 👎 to provide feedback.

…y nb_neighbors(1) searchImpl_ (the standard GpuIndex::search path) selected the filtered kernel whenever a bitset was set via sticky setSearchParams but never uploaded it, dereferencing an unallocated device buffer. Upload the bitset before gpu_hnsw_search, mirroring searchHost/searchHostInt8. extract_hnsw_layers read nb_neighbors(1) unconditionally; for a degenerate layer-0-only index (max_level==0) that indexes cum_nneighbor_per_level[2] out of bounds. Only read it when an upper layer exists (maxM is unused otherwise). Addresses Devin Review findings on PR #10. Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

premal · 2026-07-21T23:10:12Z

+                staging_ids[slot] = nbr;
+                staging_dists[slot] = dist;
+            }
+        }


Magnitude Review — staging overflow drops neighbors but marks them visited

Valid tuning note. bitmap_visit marks a node visited before it enters staging; if meta[1] (staging count) exceeds max_staging, the overflow candidate is silently dropped but remains marked visited — it can never be reconsidered.

Not a crash, and recall gates pass at current ef/search_width configs. But worth noting for future tuning: if search_width * max_degree0 newly-visited neighbors per iteration becomes common (dense graph, low ef), this could slightly reduce recall. The graceful degrade is the BF fallback path, never a silent crash.

Agreed — valid tuning note, not a correctness bug. max_staging = padded_staging_capacity(search_width, max_degree0) is a power-of-two ≥ search_width * max_degree0, so overflow only occurs when a single iteration newly-visits more than that padded bound; dropped-but-visited nodes can then not be reconsidered. Recall gates pass at current ef/search_width, and the BF fallback is the graceful degrade. Leaving as-is; noted for future tuning if we move to denser graphs / lower ef.

devin-ai-integration · 2026-07-22T04:42:34Z

Superseded by #11 — clean single-commit re-creation of this feature on branch gpu-hnsw (off current main), dropping the stacked history (OCQ-beam prototype + its removal). Final tree is identical to this branch. Continuing review on #11.

devin-ai-integration Bot and others added 29 commits July 9, 2026 06:30

fix(gpu_hnsw): bind device in ensure_filter before cudaMalloc (multi-…

12ef512

…GPU) Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

Merge pull request #7 from 6si/gpu-hnsw-filtered-search

9765576

feat(gpu_hnsw): CPU-parity filtered search (deletes/TTL/partition bitset)

Merge pull request #8 from 6si/devin/gpu-hnsw-h1-warp-coalesced

f22116e

perf(gpu_hnsw): warp-cooperative layer-0 distance (coalesced dataset loads)

Merge pull request #9 from 6si/devin/gpu-hnsw-oom-chunk

4523bda

fix(gpu_hnsw): bound visited-bitmap VRAM via nq-chunked layer-0 search

test(gpu_hnsw): add main() to TestGpuIndexHNSW (fix link error)

03cc560

Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

test(gpu_hnsw): feed int8-range data to SQ_Int8_L2 (direct-signed codec)

786568a

Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

premal self-assigned this Jul 21, 2026

devin-ai-integration Bot changed the base branch from gpu-hnsw-faiss to main July 21, 2026 13:35

devin-ai-integration Bot changed the title ~~feat(gpu_hnsw): faiss-native GpuIndexHNSW (vanilla IndexHNSW, cloner-bridged)~~ feat(gpu): GPU HNSW for faiss — GpuIndexHNSW (int8/DP4A, fp16/bf16, filtered, cloner-bridged) Jul 21, 2026

devin-ai-integration Bot commented Jul 21, 2026

View reviewed changes

Comment thread faiss/gpu/GpuIndexHNSW.cu

devin-ai-integration Bot commented Jul 21, 2026

View reviewed changes

Comment thread faiss/gpu/impl/GpuHnswBuildCommon.cuh Outdated

devin-ai-integration Bot commented Jul 21, 2026

View reviewed changes

Comment thread faiss/gpu/GpuIndexHNSW.cu

devin-ai-integration Bot commented Jul 21, 2026

View reviewed changes

premal reviewed Jul 21, 2026

View reviewed changes

Comment thread faiss/gpu/GpuIndexHNSW.cu

premal reviewed Jul 21, 2026

View reviewed changes

devin-ai-integration Bot closed this Jul 22, 2026

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

feat(gpu): GPU HNSW for faiss — GpuIndexHNSW (int8/DP4A, fp16/bf16, filtered, cloner-bridged) - #10

feat(gpu): GPU HNSW for faiss — GpuIndexHNSW (int8/DP4A, fp16/bf16, filtered, cloner-bridged)#10
devin-ai-integration[bot] wants to merge 30 commits into
mainfrom
gpu-hnsw-faiss-native

devin-ai-integration Bot commented Jul 21, 2026 •

edited

Loading

Uh oh!

devin-ai-integration Bot commented Jul 21, 2026

Uh oh!

devin-ai-integration Bot left a comment

Uh oh!

Uh oh!

Uh oh!

devin-ai-integration Bot Jul 21, 2026

Uh oh!

Uh oh!

devin-ai-integration Bot Jul 21, 2026

Uh oh!

Uh oh!

premal Jul 21, 2026

Uh oh!

devin-ai-integration Bot Jul 21, 2026

Uh oh!

devin-ai-integration Bot commented Jul 22, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Uh oh!

Conversation

devin-ai-integration Bot commented Jul 21, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Index & cloner

Search kernel (perf + correctness)

int8 accuracy (paired with the knowhere consumer)

Tests (TestGpuIndexHNSW.cpp)

Validation (CUDA arch 89, GPU CI helper)

Uh oh!

devin-ai-integration Bot commented Jul 21, 2026

🤖 Devin AI Engineer

Uh oh!

devin-ai-integration Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

devin-ai-integration Bot Jul 21, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

devin-ai-integration Bot Jul 21, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

premal Jul 21, 2026

Choose a reason for hiding this comment

Magnitude Review — staging overflow drops neighbors but marks them visited

Uh oh!

devin-ai-integration Bot Jul 21, 2026

Choose a reason for hiding this comment

Uh oh!

devin-ai-integration Bot commented Jul 22, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

devin-ai-integration Bot commented Jul 21, 2026 •

edited

Loading

Tests (`TestGpuIndexHNSW.cpp`)