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 |
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 |
CagraIndexImpl.search built the filter bitmap through Util.concatenate, which treats a BitSet with no bits set like a missing filter and sets every bit. A prefilter that excludes all rows therefore returned the excluded rows. CAGRA takes a single filter for all queries, so use its bits directly. The re-enabled CagraRandomizedIT hit this on tiny datasets whose random prefilter happened to exclude every row.
The random donor selection took candidate points from a shared atomic counter, so which receiver got which donor depended on the order in which the warps ran. Derive each receiver's candidates from its index instead, with a stride that is coprime with the number of rows, so that every receiver still finds a donor whenever one exists. The pseudo-random donor offset rotated with a function-local static counter, so a fit depended on how many fits had run before it in the same process. Rotate it with the balancing step instead.
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 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