Repository navigation
Conversation
Sum the rows and weights of each cluster in a fixed order instead of with atomics, and use cub::DeviceReduce for the centroid shift and the weight sum. Identical fits now return bitwise-identical centroids, so the convergence check no longer stops at a different iteration from run to run.
The expanded L2 distance epilogues clamp round-off of self-distances to zero when the two norms are equal and the squared distance is below an absolute tolerance. For data with a small magnitude, this also zeroes the distances between distinct points with equal norms. Spectral embeddings of graphs with several connected components are such data, which made k-means (and spectral clustering) assign points to the wrong clusters on GPUs that use the CUTLASS fused 1-NN kernel. Scale the tolerance with the norms when they are smaller than one.
k-means can leave lists empty, for example when the dataset has fewer distinct rows than n_lists. Building from device data, or from host data that fits on the device, then launched the quantization kernels with zero blocks and failed with cudaErrorInvalidValue. The streaming build skipped empty lists, which left their rotated centroids uninitialized. Compute only the rotated centroid for an empty list and skip the kernels that have no rows to process. Add a regression test with fewer distinct rows than lists.
With one row, check_graph_degree reduced both the intermediate and the final graph degree to 0, and the kNN graph build (NN-descent or IVF-PQ) then hit an illegal memory access, which aborted the process from a stream synchronization in a destructor. Throw a clear error up front instead, and keep HnswRandomizedIT from generating single-row datasets.
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
Contributor
Author
|
/ok to test ed24122 |
1 of 12 tasks
Single-partition CAGRA search computes the filtering rate from the filter bitset when the caller does not set one, and MULTI_CTA enlarges itopk_size for selective filters accordingly. The multi-partition search skipped this, so a selective per-partition filter searched with the unadjusted itopk_size and could return fewer than k neighbors although every partition held more than k rows passing the filter. cuvs-lucene then routes the query back to per-segment search, which made TestMultiSegmentGPUFilterConcurrency flaky. Derive the rate from the per-partition bitsets, sizing the shared plan for the most selective partition that accepts any row. Also limit topK in TestCuVSGaps.testVectorSearchWithAlternatingDocuments to the number of documents that have vectors.
Contributor
Author
|
/ok to test bdc80d8 |
dataset_descriptor_host::state::get() skipped waiting for the descriptor initialization event when the caller's stream handle equaled the stream that ran the initialization. The default raft stream is cudaStreamPerThread, whose handle is identical in every host thread but refers to a different stream in each. When several threads search with the same raft::resources, a thread could pick up the cached descriptor and launch a search kernel before the initialization kernel queued on another thread's stream had run, reading an uninitialized descriptor (cudaErrorMisalignedAddress / cudaErrorIllegalAddress in NEIGHBORS_MG_TEST). Always wait on the ready event.
dataset_descriptor_init_with_cache() did the LRU get and set under separate locks. Threads that missed concurrently each created a descriptor, and the later insertion replaced the earlier one while other threads could already be using it. The last owner of the replaced descriptor then freed its device memory with cudaFreeAsync on its own stream, while a search kernel on another thread's stream could still be reading it. Serialize the lookup and insertion so that concurrent searches on one raft::resources share a single descriptor.
test_build_precomputed compared the neighbor indices of two separately built indexes exactly. Both indexes hold the same PQ codes, and their distances are bitwise equal, but extend fills each list in a nondeterministic order and the search breaks distance ties by that order. Compare the distances exactly and the indices of each group of equal distances as a set, searching one extra neighbor to detect a tie across the k-th position. test_ivf_pq used the default pq_dim (8 for 10 columns). With it the inner-product recall is about 0.74 and varies by about 0.01 between builds of the same data, so recall > 0.7 failed occasionally. Use one dimension per PQ subspace, which gives about 0.94.
Contributor
Author
|
/ok to test a6b8096 |
Contributor
Author
|
/ok to test 7f16733 |
Contributor
Author
|
/ok to test fb1e1b4 |
Contributor
Author
|
/ok to test 1bf8ec1 |
Use the environment-based cub::DeviceReduce with an explicit cuda::execution::determinism::run_to_run requirement for the cluster cost, the weight sum and the centroid shift. Temporary storage comes from the caller's workspace buffer through a small memory resource, so the reductions do not allocate on every call. Also let the per-cluster row sums take any data iterator and optional weights, and process the rows in passes of at most 2^22 to bound the workspace.
Contributor
Author
|
/ok to test aaa3013 |
The test required every hit of a single query to be among the closest 3 * topK vectors. The search is approximate, and on this test's data (uniform random vectors with up to 2048 dimensions) a top-1 to top-3 search occasionally returns a valid neighbor just outside that window. Check all queries instead, allow 2 such misses in total, and keep hard checks that every result is a distinct document with a vector, scored by its own vector, in score order, with min(topK, vectors) results. Also raise TOP_K_LIMIT to 128: the search does not fail beyond 64.
Contributor
Author
|
/ok to test 4b95c30 |
Define cuvs::neighbors::detail::is_device_accessible in src/neighbors/detail/pointer_attributes.cpp instead of the header-only cuvs/util/pointer_attributes.hpp. common.hpp declares it next to expect_device_accessible_data_handle, its only user in the public headers, and the per-device ATS attribute cache is now internal to libcuvs.
Contributor
Author
|
/ok to test 5152075 |
The per-cluster codebook training listed each cluster's rows in the order given by an atomic counter, so the per-cluster training subsets depended on scheduling. Group the row indices with a stable sort by label instead, so each cluster's rows are listed in increasing order.
Do not merge. Every test job runs regardless of changed files, the conda C++, conda Python, and wheel test jobs use the union of the pull-request and nightly test matrices, and every test job needs a trivial anchor job so that re-running the anchor re-runs all test jobs without rebuilding.
Contributor
Author
|
/ok to test dbeeef4 |
Contributor
Author
|
/ok to test 23d182e |
Contributor
Author
|
Closing the combined stress run. It finished on 2026-10-12: 3,025 test-job executions (25 passes over 121 jobs), with no recurrence of any fixed flake. Summary: #2783 (comment) |
This was referenced Oct 12, 2026
rapids-bot Bot
pushed a commit
that referenced
this pull request
Oct 13, 2026
…ndTrip (#2831) `TestAcceleratedHNSWMultiLayerRoundTrip.testThreeLevelGraphCanBeReopenedAndSearched` queried with `vectors[graph.entryNode()]` and checked the hit's stored `id` against the same number. `entryNode()` is a vector ordinal, which equals the docID in the merged segment, but `vectors[]` is indexed by insertion id. The two only match while `forceMerge(1)` keeps documents in insertion order. `newIndexWriterConfig()` picks a random merge policy, and `TieredMergePolicy` or `MockRandomMergePolicy` can reorder documents. The query is then an arbitrary document's vector, and the test effectively asserts that an approximate 1-bit BQ HNSW search at k = ef = 10 finds that document. A miss there is legitimate. The test now looks up the entry node's stored id, queries with `vectors[queryId]`, and requires `hit.doc == entryNode`. The assertion stays strict for all three formats. Supporting analysis (CPU only): - A replay of the failing seed picks `TieredMergePolicy`, which reorders 232 of 256 documents. - Under the Lucene102 OSQ quantization of the test's exact 256x129 vectors, each document's own quantized score is strictly the maximum. The smallest margin is 0.39, so this is not a quantization tie. - In an emulation of the BQ graph search, an arbitrary target is missed in about 22% of trials, and the entry node's own vector in 0 of 20,000. CI evidence (reproducer #2827; 4 lucene test jobs per run: 2x 13.3.0 rtxpro6000, 12.9.2 rtxpro6000, 12.9.2 l4; each loops this test with `-Dtests.iters=100` and random seeds for 45 minutes): - Before, [run 38142877368](https://github.com/NVIDIA/cuvs/actions/runs/38142877368) at `37efbf2a`: each iteration runs both the original and the fixed query on the same index. Over 39,320 iterations per format, the original query missed 112 times for BQ, 5 for SQ and 0 for float. All of those misses were on the 676 indexes that `forceMerge(1)` reordered (`TieredMergePolicy` or `MockRandomMergePolicy`). There were 0 misses when docID == insertion id. The fixed query missed 0 times. The stress-run seed `73F002E9239F2329` reorders 232/256 docs on the GPU, and the original query missed in 6 of 20 replays. - After, [run 38149047284](https://github.com/NVIDIA/cuvs/actions/runs/38149047284) at `81f14dce`, which runs this PR's exact test file: 93,300 test iterations (31,100 per format) with 0 failures and 0 errors, and the failing seed passed 20 of 20 replays. Found by the stress run (#2793). Part of #2783. Authors: - Bradley Dice (https://github.com/bdice) - Igor Motov (https://github.com/imotov) Approvers: - Igor Motov (https://github.com/imotov) URL: #2831
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Do not merge. Combined stress run for #2783 (workstream S): this branch merges every flaky-test fix marked
FIX READYon #2783 and runs all test jobs on the union of the pull-request and nightly test matrices, then re-runs the test jobs repeatedly.The
CI stress testingcommit changespr.yamlfor this PR only (all test jobs always run,matrix_type: "pull-request,nightly", and astress-rerun-anchorjob so that the test jobs can be re-run without rebuilding).Included fixes
bdice:kmeans-deterministic-centroid-sumsbdice:kmeans-balanced-deterministicbdice:fix-l2-exp-self-distance-clampbdice:lucene-search-merge-opened-readersbdice:fix-cagra-randomized-recallbdice:fix-ivf-pq-python-flaky-testsbdice:wsC-fix-onlybdice:fix-cagra-descriptor-stream-racebdice:fix-ivf-rabitq-empty-listsbdice:fix-ivf-rabitq-exact-match-search-onlybdice:fix-cutile-bin2cbdice:fix-nntest-skipbdice:fix-lucene-randomized-hnsw-recallbdice:fix-ivf-flat-cosine-zero-norm-centersbdice:fix-hmm-pageable-memory-residencybdice:fix-lucene-multilayer-entry-query