Skip to content

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

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

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

Conversation

@devin-ai-integration

Copy link
Copy Markdown

Summary

Removes a dead ~60-line block from the vendored thirdparty/faiss/faiss/gpu/impl/GpuHnswSearchKernel.cuh that has zero callers. Mirrors the same cleanup in 6si/faiss#2.

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 under thirdparty/faiss/faiss/gpu/:

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

No functional change — pure dead-code removal. Targets gpu-hnsw-faiss (not a rewrite of that branch).

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

premal and others added 30 commits June 26, 2026 21:47
Port GPU HNSW search from raw kernel interface to FAISS GPU module:

- Add faiss::gpu::GpuIndexHNSW to vendored FAISS (thirdparty/faiss/)
  with OCQ beam search kernel, warp-cooperative distance computation,
  and support for float32/int8 data with L2/IP/cosine metrics
- GpuHnswIndexNode now wraps faiss::gpu::GpuIndexHNSW instead of
  calling raw CUDA kernel interface directly
- Register GPU_HNSW index type for fp32 and int8 data types
- Update FAISS GPU CMakeLists.txt with new source files

The kernel code (Phase 1 upper-layer greedy search, Phase 2 OCQ
beam search) is functionally identical to the previous implementation
but now lives in FAISS's gpu/impl/ directory following the
GpuIndexCagra pattern.

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>
Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
…nce metrics

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
- cmake/libs/libfaiss.cmake: add faiss_gpu_hnsw OBJECT library with GPU HNSW
  CUDA sources (conditional on WITH_CUVS) so symbols are compiled into
  libknowhere.so instead of remaining undefined
- thirdparty/faiss/faiss/gpu/GpuIndexHNSW.cu: add missing DeviceUtils.h include
  for DeviceScope class

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>
… GPU HNSW

Sync fixes from 6si/faiss 1bed5ad1:
- searchImpl_: use D2D copies for queries/distances, H2D for labels
- Stagnation counter: only count when rc >= ef to prevent premature termination

Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
The eager GPU upload in GpuHnswIndexNode::Deserialize was reading
the metric type from the config parameter. When Deserialize is called
without an explicit metric_type in the json (e.g. from the test suite
with empty config), it defaults to L2. This causes:
- IP recall 0.081 (GPU computes L2 instead of IP)
- COSINE recall 0.000 (vectors not normalized, L2 computed)

Fix: detect metric from the deserialized FAISS index type itself:
- HasInverseL2Norms dynamic_cast detects cosine indexes
- metric_type == METRIC_INNER_PRODUCT detects IP indexes

Also fix GpuIndexHNSW::copyFrom() to use the same detection.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
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>
Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Add direct search parameter passing mechanism via setSearchParams()
that stores params on the GpuIndexHNSW object. GpuHnswIndexNode::Search
now calls setSearchParams() before search() to ensure ef reaches the
kernel without depending on dynamic_cast across library boundaries.

searchImpl_() checks for direct params first, then falls back to
dynamic_cast from SearchParameters.

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

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

Add glog-based logging (LOG_KNOWHERE_INFO_) at key points in
GpuHnswIndexNode::Search to trace whether the function is called,
what ef value is used, and whether gpu_index_ is initialized.

Also add fflush(stderr) after every fprintf in CUDA code to ensure
diagnostic output is not lost in container stderr buffering.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
GpuIndex::search_ex creates temporary device allocations via
toDeviceTemporary and passes device pointers to searchImpl_.
Our searchImpl_ then does its own D2D copies internally, creating
a fragile GPU→host→GPU→host round-trip for labels that causes
SIGSEGV when the GpuIndex temp memory is freed/reused.

Add searchHost() method that takes host pointers directly:
- H2D copy queries to scratch buffer
- Run kernel
- Sync stream
- D2H copy distances and labels directly to host output
- No intermediate device temporaries, no pointer confusion

GpuHnswIndexNode::Search now calls searchHost() instead of
gpu_index_->search(), eliminating the crash.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
get_xb() can return a device pointer in GPU querynode context
(Knowhere's IndexFlat may use GPU-resident storage). Constructing
std::vector from this pointer dereferences device memory on CPU,
causing SIGSEGV (SI_CODE:2, SEGV_ACCERR).

reconstruct_n() always writes to the provided host buffer regardless
of where the underlying storage lives.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
dataset->GetTensor() returns a device pointer in GPU querynode context.
Use cudaPointerGetAttributes to detect device pointers and copy to host
before COSINE normalization and searchHost(). Also use cudaMemcpyDefault
in searchHost() for defense-in-depth.

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

cudaPointerGetAttributes requires an active CUDA context. Milvus querynode
runs Search on Folly threadpool worker threads that may not have a CUDA
context initialized, causing SIGSEGV. Query data from GetTensor() is always
plain CPU memory (aligned_vector<char>), so the device pointer check was
unnecessary. searchHost() retains cudaMemcpyDefault for safety.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Folly CPUThreadPoolExecutor worker threads may not have an active CUDA
context. DeviceScope calls cudaGetDevice which crashes without one.
Add explicit cudaSetDevice as first CUDA call to initialize context.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
StandardGpuResources pre-allocates 1.5 GiB temp buffer per instance.
With 54 segments, this totals 81 GiB exceeding L40S 96 GiB VRAM.
HNSW does not use this IVF-oriented temp buffer — set to 0.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
After copyFromWithMetric copies vectors and graph to GPU, the CPU
indexes[0] is a redundant copy. At 10M vectors with 54 segments,
keeping both copies exhausts the 62 GiB RAM on g7e.2xlarge.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Search() is const but needs to release the CPU index after GPU upload
to prevent RAM OOM. gpu_resources_ and gpu_index_ are already mutable;
indexes is not. const_cast is safe here because the mutation is
logically const (freeing a cache after populating the GPU equivalent).

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Each StandardGpuResources instance allocates 256 MiB pinned host memory
(cudaHostAlloc), a cuBLAS handle, and 4 CUDA streams. With 109 segments
per querynode, per-segment allocation wastes ~28 GiB of pinned memory
plus cuBLAS/stream overhead. The gpu-hnsw-sq branch avoids this entirely
by using raw cudaMalloc.

Fix: introduce a process-global singleton initialized once with
setTempMemory(0) and setPinnedMemory(0). HNSW search manages its own
device memory and does not need FAISS's async copy buffer or IVF temp
workspace.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
GPU HNSW stores all data in VRAM and frees the CPU copy after upload.
Without this override, the Milvus memory predictor falls back to
memoryCost=file_size (79 MB/segment), creating 48.9 GiB of phantom CPU
reservations for 616 segments even though actual RSS is only 10.4 GiB.
This causes the predictor to refuse loading more segments and eventually
OOMKill the pod.

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

StandardGpuResourcesImpl::initializeForDevice is not thread-safe —
concurrent GpuIndexHNSW constructors race on the allocs_ map assertion.
Add a process-global mutex around the GpuIndexHNSW constructor call at
both Deserialize and Search lazy-load sites. The slow copyFromWithMetric
upload runs outside the lock.

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

Fix 19: Remove fprintf+fflush from gpu_hnsw_search() and searchImpl_(),
downgrade LOG_KNOWHERE_INFO_ to LOG_KNOWHERE_DEBUG_ in Search() method.
These were called on every search request causing blocking I/O overhead.

Fix 20: Create per-segment cudaStreamNonBlocking streams instead of using
the shared StandardGpuResources default stream. With 77 segments per
querynode all serialized on one stream, GPU kernels could not overlap.
Per-segment streams allow concurrent kernel execution across segments.

Fix 21: Change max_iter formula from (ef+overflow_ef)/sw+20 to 2*ef/sw+10
to match gpu-hnsw-sq branch, reducing iterations from 170 to 110 at ef=200.

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

The overflow queue stores evicted candidates from the top-ef result set in global
memory. Each iteration, overflow_insert does insertion sort into this buffer.
With overflow_factor=2, the buffer holds 2*ef=400 entries per query. Reducing to 1
halves this to ef=200 entries, cutting global memory traffic and insertion sort
overhead per iteration. gpu-hnsw-sq has no overflow queue at all.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
select_threads_per_dist(384) was returning 4, giving 32 concurrent distances
per block. gpu-hnsw-sq uses 1 thread per distance = 128 concurrent distances.
For int8 dim=384 (384 bytes/vector), single-thread distance fits in L1 cache
and the 4x concurrency gain outweighs cooperative coalescing benefits.

This matches the gpu-hnsw-sq kernel's approach which achieved 50x speedup.

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

overflow_factor=0 by default: no overflow buffer allocation, no global memory
insertion sort per iteration, no overflow candidate re-expansion.

Remove stagnation detection (meta[3] >= 4 early exit): run for full
max_iterations like gpu-hnsw-sq. Reduce meta from 4 to 3 ints.

When overflow_ef > 0 is passed, overflow still works (backward compatible).

These changes make the kernel structurally identical to gpu-hnsw-sq's
beam search: simple sorted result buffer in shared memory, evicted
candidates dropped, no global memory overflow tracking.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
Replace byte-by-byte load_elem() with char4 vectorized loads in int8_t
distance specializations. Reads 4 bytes per load instruction instead of 1,
reducing load instruction count 4x for int8 data.

Remove coop_*_distance wrapper overhead: since threads_per_dist=1, the
cooperative machinery (chunk calc, group_mask, shuffle reduction) is
unnecessary. Call thread_*_distance directly from the kernel — each thread
independently computes one distance, matching gpu-hnsw-sq's pattern exactly.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
v20 regressed to 3,427 vec/s (2.1x slower than v19's 7,176 vec/s).
Root cause: overflow_factor=0 removes the overflow queue which acts
as the convergence mechanism — without it, newly inserted candidates
continuously create unexpanded entries, preventing early termination.

Revert overflow_factor from 0 to 1. Keep Fix 25 (char4 vectorized
int8 loads + direct thread_*_distance) which was not in v20.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
v21 eval showed -13% throughput (6,270 vs v19's 7,176 vec/s) and
-1.0pp R@1 (0.893 vs 0.903). Per-element __ldg loads via load_elem()
outperform char4 vectorized loads, likely due to register pressure
or L1 cache behavior differences.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
premal and others added 9 commits June 30, 2026 06:31
…, loop unroll

Fix 27: Multiple search performance optimizations:
- Set overflow_factor=0 (stagnation break at num_parents==0 is
  independent, so this is safe unlike Fix 24 which also removed
  stagnation detection)
- Remove __ldg from graph and inv_norms loads in layer0 and upper
  layer kernels — use L1 cache instead of texture cache for random
  access patterns (matches gpu-hnsw-sq behavior)
- Add #pragma unroll 8 to distance computation loops for ILP
- Use cudaMemcpyAsync for D2H copies with single cudaStreamSynchronize
  at the end instead of sync-then-copy pattern

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

Fix 28: Two knowhere-level search optimizations:
- Remove 5x LOG_KNOWHERE_DEBUG_ calls from Search() hot path
  (debug logging still has overhead even when disabled)
- Add double-checked locking for gpu_index_ init: skip mutex
  acquisition when gpu_index_ is already set (common case after
  Deserialize)

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
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.

At high concurrency (w>=16), the scratch_mutex was the primary
bottleneck — with ~13 segments per node, concurrent search batches
queued on each segment's mutex. The pool eliminates this serialization.

Pool uses RAII ScratchPoolGuard for exception-safe acquire/release
and condition_variable for blocking when all slots are in use.

Signed-off-by: premal <premal@6sense.com>
Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>
gpu_hnsw_search() was still reading idx.scratch which no longer exists
after the scratch pool change. Now takes GpuHnswSearchScratch& as a
parameter, and callers pass the scratch from their pool slot.

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

Sync faiss GPU HNSW kernel optimizations from 6si/faiss@e8c05ec2:

1. dp4a INT8 dot product: quantize query to int8 in shared memory,
   use __dp4a intrinsic for 4x fewer instructions (96 vs 384 ops).
2. Shared memory query cache: load query once, reuse across all
   beam search iterations.

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>
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 (mirrors 6si/faiss cleanup).

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

@devin-ai-integration

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

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant