feat(gpu_hnsw): faiss-native GpuIndexHNSW (vanilla IndexHNSW, cloner-bridged) by devin-ai-integration[bot] · Pull Request #11 · 6si/faiss

devin-ai-integration · 2026-07-22T04:39:41Z

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 is a clean single-commit re-creation of the feature off the current main (replaces the earlier stacked gpu-hnsw-faiss-native branch, whose history carried the OCQ-beam prototype and its removal). Final tree is identical to the reviewed branch.

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

Unified layer-0 kernel with native int8 DP4A 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.
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.

int8 accuracy

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. TestGpuIndexHNSW 9/9 pass on CUDA arch 89.

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

…bridged) GPU HNSW as a first-class GPU index built on vanilla faiss::IndexHNSW (Flat/SQ storage + faiss::HNSW graph), produced by the standard cloner like GpuIndexFlat/GpuIndexIVF* (CPU IndexHNSW --cpu_to_gpu--> GpuIndexHNSW, search-only; copyTo throws). Index & cloner: copyFrom maps faiss::HNSW (CSR neighbors, entry_point, levels, cum_nneighbor_per_level) to a flat device graph + uploads storage; cosine = normalize + METRIC_INNER_PRODUCT. Cloner routes faiss::IndexHNSW (excluding IndexHNSWCagra) to GpuIndexHNSW; SWIG downcast + AutoTune efSearch sweep. Search kernel: unified layer-0 kernel with native int8 DP4A (QT_8bit_direct_signed) and warp-cooperative coalesced loads; native FP16/BF16 device storage; parallel bitonic-sort + merge-path merge; CPU-parity filtered search (deletes/TTL/partition bitset) on-device with per-GPU device binding; nq-chunked layer-0 search bounds the visited-bitmap VRAM with OOB guards. int8 accuracy: INT8 L2/IP use QT_8bit_direct_signed + DP4A; knowhere re-encodes int8-cosine as fp16 (SQ_Fp16_Cosine is the representative gate). Tests (TestGpuIndexHNSW.cpp): Flat L2/IP/cosine, SQ int8 L2, SQ fp16 L2, SQ fp16/bf16 cosine, cloner dynamic type, valid IDs, distance sign/order, recall, unsupported copyTo/metric. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

devin-ai-integration · 2026-07-22T04:39:44Z

🤖 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 8 potential issues.

View 2 additional findings in Devin Review.

devin-ai-integration · 2026-07-22T04:51:37Z

+            if (sc.d_queries_i8 != nullptr && idx.use_ip && (dim % 4 == 0)) {
+                launch_kernels.template operator()<int8_t, int8_t, true>(
+                        static_cast<const int8_t*>(idx.d_dataset),
+                        idx.d_inv_norms,
+                        sc.d_queries_i8);


🟡 Reused int8 query buffer can make a float search return wrong results

The native int8 fast path is chosen whenever a leftover int8 query buffer exists on the reused scratch slot (sc.d_queries_i8 != nullptr at faiss/gpu/impl/GpuHnswSearch.cuh:450) instead of when int8 queries were actually staged this search, so a float-query search inheriting that slot silently scores against stale int8 query data.
Impact: On an inner-product int8 index, a normal search that follows an int8 search on the same index can return incorrect nearest neighbors with no error.

Stale scratch buffer + wrong path selection mechanism

GpuHnswSearchScratch::d_queries_i8 is allocated only when ensure(..., use_i8_queries=true) is called, which happens exclusively in searchHostInt8 (faiss/gpu/GpuIndexHNSW.cu:363). The buffer is grow-only: it is never freed or nulled between searches (faiss/gpu/impl/GpuHnswTypes.cu:98-105). Slots are pooled and reused across all search entry points of the same GpuIndexHNSW (they share deviceIndex_->scratch_pool).

In gpu_hnsw_search the INT8 dataset branch selects the DP4A path via if (sc.d_queries_i8 != nullptr && idx.use_ip && (dim % 4 == 0)) (faiss/gpu/impl/GpuHnswSearch.cuh:450). The float paths (searchImpl_, searchHost) only write sc.d_queries and never populate sc.d_queries_i8. If a prior searchHostInt8 allocated d_queries_i8 on the acquired slot, a subsequent float search on an int8/IP index takes the DP4A branch and reads the stale int8 query bytes rather than the freshly-uploaded float queries.

A robust fix would gate the DP4A path on a per-search flag (e.g. set on the scratch slot when int8 queries are actually uploaded this search, cleared by the float paths) rather than on mere buffer existence.

Prompt for agents

The DP4A int8 fast path in gpu_hnsw_search (faiss/gpu/impl/GpuHnswSearch.cuh:450) is selected whenever sc.d_queries_i8 != nullptr. However d_queries_i8 is a grow-only buffer on a pooled scratch slot (GpuHnswSearchScratch, allocated only by ensure(..., use_i8_queries=true) which is called only from searchHostInt8). Because slots are reused across all search entry points of the same GpuIndexHNSW, a float-query search (searchImpl_ or searchHost) that acquires a slot previously used by searchHostInt8 will see a stale non-null d_queries_i8 and take the DP4A path, reading stale int8 query bytes instead of the float queries it just uploaded to d_queries. This yields silently wrong results on int8/inner-product indexes. Fix by gating the DP4A path on an explicit per-search signal (e.g. a boolean on the scratch slot set true when int8 queries are uploaded this search and reset to false by the float upload paths), rather than on buffer existence.

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

devin-ai-integration · 2026-07-22T04:51:38Z

+    int dim = static_cast<int>(idx.dim);
+    int64_t nelem = static_cast<int64_t>(nq) * dim;
+
+    sc.ensure(nq, k, dim, static_cast<int>(idx.n_rows), /*use_i8_queries=*/true);


🟡 Source lines exceed the repository's 80-character limit

Two new lines are wider than the project's required maximum line width (faiss/gpu/GpuIndexHNSW.cu:340 and faiss/gpu/GpuIndexHNSW.cu:363, both 81 characters), breaking the coding-style rule that all lines stay within 80 characters.
Impact: The change does not conform to the project's documented style, so it can fail lint/style checks.

CONTRIBUTING.md 80-char rule

CONTRIBUTING.md states "80 character line length (both for C++ and Python)". Lines 340 and 363 in faiss/gpu/GpuIndexHNSW.cu are 81 characters.

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

devin-ai-integration · 2026-07-22T04:51:39Z

+    int64_t dim = idx.dim;
+    // Flat cosine (stored_inv_norms == nullptr): the stored vectors are the
+    // exact originals, so normalizing them in place yields cosine via plain
+    // inner product. Lossy-SQ cosine decoded to fp32 (stored_inv_norms != null):


🟡 Comment line exceeds the repository's 80-character limit

A new comment line is wider than the project's required maximum line width (faiss/gpu/impl/GpuHnswBuildCommon.cuh:261, 81 characters), breaking the coding-style rule that all lines stay within 80 characters.
Impact: The change does not conform to the project's documented style, so it can fail lint/style checks.

CONTRIBUTING.md 80-char rule

CONTRIBUTING.md states "80 character line length (both for C++ and Python)". Line 261 in faiss/gpu/impl/GpuHnswBuildCommon.cuh is 81 characters.

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

devin-ai-integration · 2026-07-22T04:51:40Z

+
+    // Layer-0 is templated on the dataset type (DataT), the layer-0 query type
+    // (QueryT: float generic, int8_t for the native DP4A path) and USE_DP4A.
+    // The upper-layer greedy descent always uses the fp32 queries (sc.d_queries).


🟡 Multiple lines exceed the repository's 80-character limit

Several new lines are wider than the project's required maximum line width (starting at faiss/gpu/impl/GpuHnswSearch.cuh:122), breaking the coding-style rule that all lines stay within 80 characters.
Impact: The change does not conform to the project's documented style, so it can fail lint/style checks.

CONTRIBUTING.md 80-char rule

CONTRIBUTING.md states "80 character line length (both for C++ and Python)". In faiss/gpu/impl/GpuHnswSearch.cuh lines 122, 132, 170, 289, and 368 exceed 80 characters.

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

devin-ai-integration · 2026-07-22T04:51:41Z

+            // pair into registers and decides whether to swap, (2) __syncthreads
+            // so all reads complete before any write, (3) threads that must swap


🟡 Multiple lines exceed the repository's 80-character limit

Several new lines are wider than the project's required maximum line width (starting at faiss/gpu/impl/GpuHnswSearchKernel.cuh:366), breaking the coding-style rule that all lines stay within 80 characters.
Impact: The change does not conform to the project's documented style, so it can fail lint/style checks.

CONTRIBUTING.md 80-char rule

CONTRIBUTING.md states "80 character line length (both for C++ and Python)". In faiss/gpu/impl/GpuHnswSearchKernel.cuh lines 366, 367, 380, 445, and 977 exceed 80 characters.

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

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

+namespace gpu {
+
+// Test-only fault injection for the device-upload path (consulted by
+// GPU_HNSW_BUILD_CUDA_CHECK in GpuHnswBuildCommon.cuh). Production code never arms it


🟡 Multiple lines exceed the repository's 80-character limit

Several new lines are wider than the project's required maximum line width (starting at faiss/gpu/impl/GpuHnswTypes.h:41, 86 characters), breaking the coding-style rule that all lines stay within 80 characters.
Impact: The change does not conform to the project's documented style, so it can fail lint/style checks.

CONTRIBUTING.md 80-char rule

CONTRIBUTING.md states "80 character line length (both for C++ and Python)". In faiss/gpu/impl/GpuHnswTypes.h lines 41, 46, 231, and 262 exceed 80 characters.

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

devin-ai-integration · 2026-07-22T04:51:44Z

+        slot->scratch.device = device_;
+        SCRATCH_CUDA_CHECK(cudaSetDevice(device_));
+        SCRATCH_CUDA_CHECK(
+                cudaStreamCreateWithFlags(&slot->stream, cudaStreamNonBlocking));


🟡 Source line exceeds the repository's 80-character limit

A new line is wider than the project's required maximum line width (faiss/gpu/impl/GpuHnswTypes.cu:184, 81 characters), breaking the coding-style rule that all lines stay within 80 characters.
Impact: The change does not conform to the project's documented style, so it can fail lint/style checks.

CONTRIBUTING.md 80-char rule

CONTRIBUTING.md states "80 character line length (both for C++ and Python)". Line 184 in faiss/gpu/impl/GpuHnswTypes.cu is 81 characters.

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

devin-ai-integration · 2026-07-22T04:51:45Z

+}
+
+// INT8 storage via QT_8bit_direct_signed (native DP4A path). dim % 4 == 0 so
+// the DP4A kernel is exercised; quantization lowers the recall bar. The input is


🟡 Comment line exceeds the repository's 80-character limit

A new comment line is wider than the project's required maximum line width (faiss/gpu/test/TestGpuIndexHNSW.cpp:179, 81 characters), breaking the coding-style rule that all lines stay within 80 characters.
Impact: The change does not conform to the project's documented style, so it can fail lint/style checks.

CONTRIBUTING.md 80-char rule

CONTRIBUTING.md states "80 character line length (both for C++ and Python)". Line 179 in faiss/gpu/test/TestGpuIndexHNSW.cpp is 81 characters.

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

Devin Review (faiss #11): - 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. Scratch slots are pooled and d_queries_i8 stays allocated after an int8 search, so a later fp32-query search on the same slot (e.g. searchHost/searchImpl_ on an int8 index) could take the DP4A path and score against stale int8 query data. Gating on the staged flag makes the fp32 fallback correct for reused slots. - Wrap comment/source lines to the 80-char CONTRIBUTING.md limit (no behavior change). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

…arch() tests searchImpl_ (the faiss-standard GpuIndex::search() device-pointer path) no longer round-trips labels through the host: a new convert_labels_kernel maps the uint64 neighbor ids to idx_t (UINT64_MAX -> -1) directly on the slot stream, writing straight into the caller's device output. Drops the host staging alloc, the D2H copy, the CPU convert loop, and the H2D copy back; a single end-of-search stream sync remains to order the private slot stream before GpuIndex::search copies outputs back. Production searchHost/searchHostInt8 are unaffected (they already avoid the round-trip). Adds faiss/gpu/test/test_gpu_index_hnsw.py exercising the SWIG-exposed search() path: index_factory(HNSW32,{Flat,SQ8,SQfp16}) -> index_cpu_to_gpu -> GpuIndexHNSW.search() with SearchParametersGpuHNSW, asserting GPU<->CPU recall parity and label/sentinel correctness for L2 and IP (cosine). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

index_factory(HNSW32,SQ8) builds an IndexHNSWSQ whose scalar quantizer must be trained before add(); the test now calls train() first (a no-op for Flat/fp16/ bf16). Fixes the test_sq8_l2 is_trained assertion. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

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

Sync gpu-hnsw with upstream facebookresearch/faiss main (34 commits). Resolved faiss/gpu/GpuCloner.cpp: keep the IndexHNSW.h include and adopt upstream's CAGRA guard (USE_NVIDIA_CUVS && !FAISS_CUVS_NO_CAGRA). Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

Apply clang-format-21 (the version faiss CI enforces) to the GPU HNSW sources/tests and add an Unreleased CHANGELOG entry referencing facebookresearch#5458. Formatting-only; no functional change. Signed-off-by: premal <premal@6sense.com> Co-Authored-By: Devin AI <158243242+devin-ai-integration[bot]@users.noreply.github.com>

…upport (facebookresearch#5288) Summary: Adds IVF-PQ (inverted file with product quantization) index support to the Metal GPU backend - Add `MetalIndexIVFPQ` with full train/add/search/reset/copyFrom/copyTo support - Add `MetalIVFPQImpl` GPU-resident IVF list storage for PQ codes (segment allocator, same pattern as IVFFlat) - Support 8-bit product quantization with precomputed per-query lookup tables - Support both L2 and inner product metrics - Residual encoding when `by_residual=true` (default for L2) - CPU-side PQ lookup table computation with precomputed tables optimization for L2 - GPU scan path via `runMetalIVFPQFullSearch` with CPU LUT fallback via `runMetalIVFPQScan` - Update `MetalCloner` to support IVFPQ in `index_cpu_to_metal_gpu` / `index_metal_gpu_to_cpu` ## Changes - **New:** `MetalIndexIVFPQ.h/.mm` - IVFPQ index class (train, add, search, reset, copyFrom/copyTo, cloner support) - **New:** `impl/MetalIVFPQ.h/.mm` - GPU-resident IVF list storage with segment allocator for PQ codes - **New:** `test/TestMetalIndexIVFPQ.mm` - 4 C++ tests (L2, IP, reset, CPU↔GPU round-trip) - **Modified:** `test/CMakeLists.txt` - added TestMetalIndexIVFPQ build target ## Differences from CUDA IVFPQ **Training:** Delegates to CPU (`IndexIVFPQ::train`). CUDA can train on GPU. Same rationale as IVFFlat - training is a one-time cost. **Add path:** Coarse quantization and PQ encoding run on CPU, then codes are copied to GPU storage. CUDA does both on GPU. On Apple Silicon with unified memory, the copy cost is minimal. **Residual encoding:** When `by_residual=true`, residuals (x - coarse_centroid) are computed on CPU before PQ encoding. CUDA computes residuals on GPU. Functionally equivalent. **Lookup tables:** PQ distance lookup tables are computed on CPU and uploaded to GPU for the scan phase. CUDA computes LUTs on GPU. CPU LUT computation is fast relative to the scan and avoids a separate GPU kernel launch. **IVF list storage:** Same segment allocator pattern as IVFFlat - single contiguous buffer rather than CUDA's per-list `DeviceVector` allocations. ## Note FP16 coarse quantizer and GPU merge kernel are planned optimizations for a future PR. Both apply across all IVF index types (IVFFlat, IVFPQ, IVFSQ). ## Build and test ```bash cmake -B build \ -DFAISS_ENABLE_GPU=OFF \ -DFAISS_ENABLE_METAL=ON \ -DBUILD_TESTING=ON \ -DCMAKE_BUILD_TYPE=Release \ -DCMAKE_PREFIX_PATH="$(brew --prefix libomp)" \ . cmake --build build --target faiss faiss_metal TestMetalIndexIVFPQ -j$(sysctl -n hw.logicalcpu) cd build && ctest -R TestMetalIndexIVFPQ --output-on-failure ``` Pull Request resolved: facebookresearch#5288 Reviewed By: alibeklfc Differential Revision: D113037852 Pulled By: mnorris11 fbshipit-source-id: 5fdb729b10ad9884a673fc3483543b4c5fcdb621

Summary: - Use `is_similarity_metric()` for IVF heap initialization, merging, reordering, and result handling. - Add Jaccard IVFFlat regression coverage across parallel modes 0-3. ## Why Jaccard scanners use similarity ordering, but `IndexIVF::search_preassigned` selected distance heaps for every metric except inner product. The resulting `FLT_MAX` threshold rejected all Jaccard candidates, returning `-1` even when the only inverted list was fully scanned. Fixes facebookresearch#5399. ## Validation - [x] `cmake --build build --target faiss_test --config Release -j 8` - [x] `ctest --test-dir build --output-on-failure -C Release` (251 tests, 5 skipped) - [x] `clang-format-21 --dry-run --Werror faiss/IndexIVF.cpp tests/test_ivf_index.cpp` ## Notes for reviewers Python and GPU suites were not run locally; the fix and regression exercise the shared C++ IVF path. Pull Request resolved: facebookresearch#5408 Reviewed By: pankajsingh88, mnorris11 Differential Revision: D112830842 Pulled By: bshethmeta fbshipit-source-id: c6431d2bee1b4fecd2e46cfda397722947a5407a

…per call (facebookresearch#5448) Summary: Pull Request resolved: facebookresearch#5448 ## Background Every `IndexHNSW::search()` call enters an `#pragma omp parallel` region in which each thread constructs its own `VisitedTable` via `VisitedTable::create(ntotal, ...)`. For the versioned-array strategy (`VisitedTableVector`), a fresh array allocation is paid **per thread, per Search() call** - this is an O(ntotal) allocation plus zero-fill overhead. It's freed every time the region exits. ## Problem When a statically-built index is searched repeatedly with small or even single-query batches, this per-call alloc+zero can dominate the actual graph traversal (the search visits only a handful of nodes, but still pays to allocate and zero the whole `ntotal`-byte array). ## Solution This adds `VisitedTable::get_reusable()` API to return a reference to the `thread_local` VisitedTable. The O(size) versioned array is **allocated once per thread and reused across all subsequent Search() calls**. This changed is applied to 3 code paths for search (`IndexHNSW.cpp`): - `hnsw_search` - `search_level_0` - only used by IndexHNSWCagra - `IndexHNSWCagra::range_search` - only used by IndexHNSWCagra Why `thread_local` rather than a member of the index: - It is inherently per-OS-thread, so it survives across `search()` calls (OpenMP reuses its worker pool) and each thread gets its own table with no cross-thread coordination. - faiss supports calling `const` search concurrently on one index from multiple threads. A member array indexed by `omp_get_thread_num()` would let two concurrent searches race on slot 0; `thread_local` cannot, since each OS thread has its own copy. - The table grows on demand (`ensure_size`) if `ntotal` increases and never shrinks. Retained memory is bounded by the array-vs-hash-set threshold (D112025912). ## Correctness `get_reusable()` calls `advance()` before returning the table. A prior search that threw exception without correctly `advance()` may have left version stamps NOT updated; without the reset those would read as spurious "visited" hits on the next search on that thread. `advance()` - might be redundant - in `get_reusable()` guarantees correct query version id. ## Minor notes The `IndexHNSW2Level` mixed-search path is deliberately left on per-search `create()` rather than `get_reusable()`. It is a legacy path that uses the tri-state visited flags of `search_from_candidates_2` (two `advance()` calls per query), whose reset semantics differ from the bi-state paths above; keeping a fresh per-search table there leaves its behavior unchanged and scopes reuse to the paths where a single `advance()` at handout is correct. Reviewed By: mnorris11, pankajsingh88 Differential Revision: D112042940 fbshipit-source-id: c6be0effdb2f52d21f6bb16db2d2ad0972a71bc0

Summary: Pull Request resolved: facebookresearch#5461 Reviewed By: limqiying Differential Revision: D112521498 fbshipit-source-id: 5ce36c751363207d9a8375c30e153d7aa5ba88e9

…#5462) Summary: Pull Request resolved: facebookresearch#5462 Reviewed By: limqiying Differential Revision: D112521364 fbshipit-source-id: 8aaa62cf8536280235b55ddf81a5e60dc2cea3c6

An empty IndexHNSW (ntotal==0) reaching the GPU cloner is now left to the "not implemented on GPU" fall-through instead of being routed to GpuIndexHNSW::copyFrom (which throws "index must not be empty"). That throw did not match the "not implemented on GPU" substring GpuIndexIVF::copyFrom keys on, so cloning an untrained IVF_HNSW (HNSW coarse quantizer, ntotal==0) to the GPU with allowCpuCoarseQuantizer=true regressed from a CPU-coarse-quantizer fallback into a hard error (TestGpuAutoTune::test_params). GpuIndexHNSW is search-only, so an empty graph has nothing to upload; populated standalone HNSW indexes still clone to the GPU 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>

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

premal self-assigned this Jul 22, 2026

devin-ai-integration Bot mentioned this pull request Jul 22, 2026

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

Closed

devin-ai-integration Bot commented Jul 22, 2026

View reviewed changes

premal and others added 5 commits July 22, 2026 06:09

docs(gpu_hnsw): document GpuIndexHNSW (CPU-built graph, GPU search) i…

ef0c907

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

This was referenced Jul 23, 2026

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

Closed

feat: GPU HNSW search with parallel beam-search kernel — native FP32/FP16/BF16/INT8 #1

Closed

premal and others added 9 commits July 23, 2026 03:27

facebook-unused-include-check in IndexIVFPQ.cpp (facebookresearch#5461)

ab3a5a9

Summary: Pull Request resolved: facebookresearch#5461 Reviewed By: limqiying Differential Revision: D112521498 fbshipit-source-id: 5ce36c751363207d9a8375c30e153d7aa5ba88e9

facebook-unused-include-check in distances_avx2.cpp (facebookresearch…

d4564bd

…#5462) Summary: Pull Request resolved: facebookresearch#5462 Reviewed By: limqiying Differential Revision: D112521364 fbshipit-source-id: 8aaa62cf8536280235b55ddf81a5e60dc2cea3c6

Merge remote-tracking branch 'upstream/main' into gpu-hnsw

fe8e9fd

bench(gpu_hnsw): SIFT1M CPU-vs-GPU HNSW benchmark (QPS + recall)

76a8d94

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

Navigation Menu

Uh oh!

feat(gpu_hnsw): faiss-native GpuIndexHNSW (vanilla IndexHNSW, cloner-bridged) - #11

feat(gpu_hnsw): faiss-native GpuIndexHNSW (vanilla IndexHNSW, cloner-bridged)#11
devin-ai-integration[bot] wants to merge 15 commits into
mainfrom
gpu-hnsw

devin-ai-integration Bot commented Jul 22, 2026

Uh oh!

devin-ai-integration Bot commented Jul 22, 2026

Uh oh!

devin-ai-integration Bot left a comment

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

		// pair into registers and decides whether to swap, (2) __syncthreads
		// so all reads complete before any write, (3) threads that must swap

Uh oh!

Conversation

devin-ai-integration Bot commented Jul 22, 2026

Summary

Index & cloner

Search kernel

int8 accuracy

Tests (TestGpuIndexHNSW.cpp)

Uh oh!

devin-ai-integration Bot commented Jul 22, 2026

🤖 Devin AI Engineer

Uh oh!

devin-ai-integration Bot left a comment

Choose a reason for hiding this comment

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Choose a reason for hiding this comment

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Choose a reason for hiding this comment

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Choose a reason for hiding this comment

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Choose a reason for hiding this comment

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Choose a reason for hiding this comment

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Choose a reason for hiding this comment

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Choose a reason for hiding this comment

Uh oh!

devin-ai-integration Bot Jul 22, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Tests (`TestGpuIndexHNSW.cpp`)