Skip to content

refactor: remove dead cooperative-distance code from GPU HNSW kernel#2

Closed
devin-ai-integration[bot] wants to merge 27 commits into
gpu-hnsw-faissfrom
devin/gpu-hnsw-rm-deadcode
Closed

refactor: remove dead cooperative-distance code from GPU HNSW kernel#2
devin-ai-integration[bot] wants to merge 27 commits into
gpu-hnsw-faissfrom
devin/gpu-hnsw-rm-deadcode

Conversation

@devin-ai-integration

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

Copy link
Copy Markdown

Summary

Removes a dead ~60-line block from faiss/gpu/impl/GpuHnswSearchKernel.cuh that has zero callers.

The GPU HNSW search kernel settled on a 1-thread-per-distance design. An earlier warp-cooperative distance path was left behind but never wired in:

// deleted — defined, never called:
__device__ int   select_threads_per_dist(int dim)   // always returned 1
__device__ float coop_l2_distance(...)              // 4-thread cooperative L2
__device__ float coop_ip_distance(...)              // 4-thread cooperative IP

Verified no references remain anywhere under faiss/gpu/:

grep -rn 'coop_l2_distance\|coop_ip_distance\|select_threads_per_dist' faiss/gpu/  → (none)

No functional change — pure dead-code removal. Targets gpu-hnsw-faiss (not a rewrite of that branch). The identical vendored copy in 6si/knowhere is cleaned up in a parallel PR.

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

premal and others added 27 commits June 26, 2026 21:41
@premal @devin-ai-integration
feat: add GpuIndexHNSW — GPU HNSW search with OCQ beam search kernel
b1ec1a9 
Port GPU HNSW kernel from Knowhere into FAISS GPU module:

- GpuIndexHNSW: public API extending GpuIndex, supports copyFrom(IndexHNSW*)
  for both IndexHNSWFlat (float32) and IndexHNSWSQ (int8) storage types
- Phase 1 upper-layer greedy search kernel (one warp per query)
- Phase 2 layer-0 beam search with Overflow Candidate Queue (OCQ)
  maintaining sorted result buffer in shared memory and secondary
  overflow queue in global memory to prevent premature pruning
- Warp-cooperative distance computation (L2, IP, cosine with inv-norms)
- GpuHnswDeviceIndex: device-side index structure with dense layer-0
  graph and sparse upper-layer representation
- Build utilities to convert faiss::IndexHNSW CSR graph to dense GPU format

Follows GpuIndexCagra pattern for FAISS GPU integration.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
style: apply clang-format to GPU HNSW files
1a17c3d 
Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: add missing DeviceUtils.h include in GpuIndexHNSW.cu
72f30f2 
Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: device pointer handling in searchImpl_ and stagnation counter bug
1bed5ad 
searchImpl_ receives device pointers from GpuIndex::search, not host:
- queries: cudaMemcpyHostToDevice -> cudaMemcpyDeviceToDevice
- distances: cudaMemcpyDeviceToHost -> cudaMemcpyDeviceToDevice
- labels: stage through host for uint64->idx_t conversion, then H2D back

Stagnation counter: only count stagnation when result list is full (rc >= ef).
Previously, when rc < ef both prev_worst and new_worst were FLT_MAX, so the
'no improvement' condition was always true, causing premature termination
after just 4 main-loop iterations regardless of search progress.

Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: use GpuResources default stream, add cudaMalloc error checks
ebd351d 
Two fixes for SIGSEGV crash at nb=100K:

1. Stream mismatch: searchImpl_ used a custom cudaStreamNonBlocking
   stream but GpuIndex::search() copies query data on the default
   stream. Non-blocking streams don't synchronize with the default
   stream, creating a race on query data reads. Now uses the
   GpuResources default stream (matching GpuIndexCagra pattern).

2. Unchecked cudaMalloc: All cudaMalloc calls in ensure() were
   unchecked. If any allocation fails, the kernel accesses null/junk
   pointers causing SIGSEGV. Added SCRATCH_CUDA_CHECK wrapper.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: add setSearchParams() to bypass dynamic_cast for ef propagation
17d5b42 
Add direct search parameter passing mechanism via setSearchParams()
that stores params on the GpuIndexHNSW object. searchImpl_() checks
for direct params first, then falls back to dynamic_cast from
SearchParameters. This ensures ef is correctly propagated to the
kernel even if dynamic_cast fails across library boundaries.

Also adds diagnostic fprintf logging to trace ef values through
searchImpl_ and gpu_hnsw_search for debugging.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: add fflush(stderr) after fprintf diagnostics in GPU HNSW
41fe7c2 
Ensure diagnostic output is not lost due to stderr buffering
in container environments. Matches knowhere vendored version.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: add searchHost() to bypass GpuIndex::search SIGSEGV
69e03d0 
Add searchHost() that takes host pointers directly, bypassing
GpuIndex::search_ex temp allocation chain. Matches knowhere.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: use reconstruct_n instead of get_xb() in from_faiss_hnsw_flat
2e1f501 
get_xb() can return a device pointer in GPU context. Matches knowhere.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: use cudaMemcpyDefault in searchHost() for device query pointer s…
b14b980 
…afety

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: cudaSetDevice before DeviceScope in searchHost for Folly threads
e23e0df 
Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: eagerly free staging buffers after cudaMemcpy to prevent RAM OOM
5152c0b 
When loading many segments in parallel, staging vectors (signed_codes,
h_layer0_flat, h_vectors, h_inv_norms) accumulated in heap because they
were only freed at function return. With 8+ querynodes loading segments
concurrently, these buffers push past the pod memory limit.

Fix: std::vector<T>().swap(v) immediately after each cudaMemcpy to
release staging memory before the next allocation.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
perf: remove search hot-path logging, per-segment CUDA streams, reduc…
3de52a0 
…e max_iter

Fix 19: Remove fprintf+fflush from gpu_hnsw_search() and searchImpl_().
Fix 20: Per-segment cudaStreamNonBlocking for concurrent kernel execution.
Fix 21: max_iter formula 2*ef/sw+10 (matching gpu-hnsw-sq branch).

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
perf: reduce overflow_factor from 2 to 1
398dfcc 
Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
perf: use 1 thread per distance instead of 4-thread cooperative
d09d338 
Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
perf: eliminate overflow queue + stagnation detection (match gpu-hnsw…
e887b1a 
…-sq)

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
perf: vectorized int8 loads (char4) + remove coop wrapper overhead
3209e2d 
Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: revert overflow_factor to 1 (Fix 26)
19bcae3 
v20 regressed to 3,427 vec/s (2.1x slower than v19). overflow_factor=0
removes the convergence mechanism. Revert to 1.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
revert: remove char4 int8 specializations (Fix 25 caused regression)
7ed80d5 
Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
perf: overflow_factor=0, remove __ldg from graph/inv_norms, async D2H…
81f1b90 
…, loop unroll

Fix 27: Multiple search performance optimizations synced from knowhere:
- Set overflow_factor=0 (stagnation break is independent)
- Remove __ldg from graph and inv_norms loads — use L1 cache
- Add #pragma unroll 8 to distance computation loops
- Use cudaMemcpyAsync for D2H with single sync at end
- Sync copyFrom/copyFromWithMetric changes from knowhere

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
perf: replace scratch_mutex with concurrent scratch pool (Fix 29)
950324d 
Replace single-scratch + mutex serialization with a pool of 4 scratch
slots, each with its own CUDA stream. This allows up to 4 concurrent
GPU searches per segment instead of serializing all searches through
one mutex.

Synced from knowhere thirdparty/faiss.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @devin-ai-integration
fix: pass scratch ref to gpu_hnsw_search (Fix 29 build fix)
122e181 
Synced from knowhere thirdparty/faiss.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
@premal @claude @devin-ai-integration
fix: lazy init scratch pool to avoid OOM on segment load
c5813c2 
CUDA streams are now created on first acquire() instead of in the
constructor. This prevents stream/context memory from inflating the
process RSS during segment loading, which caused Milvus memory
estimator to predict 114 GB and refuse to load segments.

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 @devin-ai-integration
perf: dp4a INT8 distance + shared memory query cache
e8c05ec 
Two kernel optimizations for INT8 HNSW search:

1. dp4a INT8 dot product: For int8 datasets, quantize query to int8
   in shared memory and use __dp4a intrinsic for distance computation.
   Reduces instruction count by ~4x (96 dp4a ops vs 384 FMA ops for
   dim=384). Query quantization error is <1% for ranking purposes.

2. Shared memory query cache: Load query vector into shared memory
   once at kernel start. All threads read from shared memory instead
   of L1/global for every distance computation across hundreds of
   beam search iterations.

Both optimizations apply to the layer0_beam_search_kernel. Float
datasets continue using the existing float-based distance functions.
Upper layer search is unchanged (single distance per warp, low impact).

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>
@devin-ai-integration
Revert "perf: dp4a INT8 distance + shared memory query cache"
54a2473 
This reverts commit e8c05ec.
@devin-ai-integration
Merge remote-tracking branch 'upstream/main' into gpu-hnsw-faiss
2708bc1
@devin-ai-integration
refactor: remove dead cooperative-distance code from GPU HNSW kernel
eabc596 
select_threads_per_dist(), coop_l2_distance() and coop_ip_distance() were
superseded by the 1-thread-per-distance path and have no callers. Remove
the dead block.

Signed-off-by: Devin AI <devin-ai-integration[bot]@users.noreply.github.com>
@devin-ai-integration

devin-ai-integration Bot commented Jul 8, 2026

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

This was referenced Jul 8, 2026
Closed
Closed
@devin-ai-integration

devin-ai-integration Bot commented Jul 9, 2026

Copy link
Copy Markdown
Author

Superseded: this content is now folded directly into the clean gpu-hnsw-faiss branch (rebased onto synced upstream master, dead-code/GPU_TIMING cleanups included). Closing.

@devin-ai-integration devin-ai-integration Bot deleted the devin/gpu-hnsw-rm-deadcode branch July 9, 2026 06:31
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

1 participant

@premal