Promote filtering_rate to base search_params; honor it in brute_force - #2120
Open
maxwbuckley wants to merge 5 commits into
Open
Promote filtering_rate to base search_params; honor it in brute_force#2120maxwbuckley wants to merge 5 commits into
maxwbuckley wants to merge 5 commits into
Conversation
This was referenced May 23, 2026
📝 WalkthroughSummary by CodeRabbit
WalkthroughThis PR introduces an optional ChangesFiltering Rate Hint Optimization
Estimated code review effort🎯 3 (Moderate) | ⏱️ ~25 minutes Suggested labels
Suggested reviewers
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (1)
cpp/tests/neighbors/brute_force_prefiltered.cu (1)
1084-1130: ⚡ Quick winAdd one high-sparsity hint case to cover the lazy-
nnzCSR branch.Current inputs top out at
sparsity = 0.5, soRunWithFilteringRateHint()never exercises the newsparsity >= 0.9path where exactnnzis computed lazily. Adding at least one> 0.9case for bitmap and bitset would close this gap.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cpp/tests/neighbors/brute_force_prefiltered.cu` around lines 1084 - 1130, Add at least one test entry to the selectk_inputs vector that triggers the lazy-nnz CSR branch by using a sparsity >= 0.9; update the PrefilteredBruteForceInputs initializer list (selectk_inputs) to include a case with high sparsity (e.g. 0.95) so RunWithFilteringRateHint() exercises the sparsity >= 0.9 path (add one for bitmap/bitset coverage if your test matrix type selection is driven by the distance/type field such as cuvs::distance::DistanceType::L2Expanded or InnerProduct).
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@cpp/src/neighbors/detail/knn_brute_force.cuh`:
- Around line 623-627: The code currently enables hint mode for any non-negative
params.filtering_rate; change this to only enable hint mode when filtering_rate
is in [0.0, 1.0) by replacing the use_hint initialization with an explicit range
check (e.g., const bool use_hint = (params.filtering_rate >= 0.0f &&
params.filtering_rate < 1.0f)); additionally add an explicit validation branch
near that declaration (in the knn_brute_force.cuh scope where use_hint and
params.filtering_rate are used) that treats values >= 1.0f as invalid—either
disable the hint and emit a clear error/exception (e.g., throw
std::invalid_argument or return an error) or log a warning before proceeding—so
out-of-contract inputs do not silently alter path selection.
---
Nitpick comments:
In `@cpp/tests/neighbors/brute_force_prefiltered.cu`:
- Around line 1084-1130: Add at least one test entry to the selectk_inputs
vector that triggers the lazy-nnz CSR branch by using a sparsity >= 0.9; update
the PrefilteredBruteForceInputs initializer list (selectk_inputs) to include a
case with high sparsity (e.g. 0.95) so RunWithFilteringRateHint() exercises the
sparsity >= 0.9 path (add one for bitmap/bitset coverage if your test matrix
type selection is driven by the distance/type field such as
cuvs::distance::DistanceType::L2Expanded or InnerProduct).
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 73fd1e69-bb28-41e6-8152-51e18206fd9f
📒 Files selected for processing (5)
cpp/include/cuvs/neighbors/cagra.hppcpp/include/cuvs/neighbors/common.hppcpp/src/neighbors/brute_force.cucpp/src/neighbors/detail/knn_brute_force.cuhcpp/tests/neighbors/brute_force_prefiltered.cu
aamijar
added
non-breaking
`filtering_rate` (a hint of the fraction of items filtered out by the sample filter) previously lived on `cagra::search_params` only. CAGRA used it to size `itopk_size` and, when the user left it at its default of `-1.0`, called `bitset_view::count(res)` on every search — a GPU popcount reduction + host sync that adds measurable latency. `brute_force` filtered search also called `bitset_view::count(res)` on every search to compute `sparsity` (used to choose between the dense tiled GEMM path and the sparse CSR path), but had no user-facing knob to skip the auto-detection. This change: - Moves `float filtering_rate = -1.0` from `cagra::search_params` to the base `cuvs::neighbors::search_params`. `cagra::search_params` inherits it; existing code that accesses `params.filtering_rate` is unaffected. - Plumbs `brute_force::search_params` through `detail::search` and `brute_force_search_filtered`. When `params.filtering_rate >= 0`, the hint is used directly as `sparsity` and the per-search popcount is skipped on the dense path. The CSR path still needs an exact non-zero count to size the matrix, so popcount runs lazily there. - Algorithms that don't use the hint (`ivf_flat`, `ivf_pq`, `hnsw`) ignore the new base field; their `search_params` inherit it but nothing reads it. Tests: extends `brute_force_prefiltered.cu` with `ResultWithFilteringRateHint` cases for both bitmap and bitset fixtures (float and half), reusing the existing parameter matrix and asserting that results match the auto-detect path. Closes NVIDIA#1960. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
maxwbuckley
force-pushed
the
filtering-rate-base-search-params
branch
from
July 14, 2026 06:07
71decfc to
d99e6a0
Compare
The base search_params docstring already defines the contract: negative means "auto-detect from the filter", and [0.0, 1.0) is trusted as a hint. The code only checked `>= 0.0f`, so an out-of-contract value (e.g. a caller passing a percentage like 50.0, or 1.5) was silently accepted as `sparsity`, forcing the sparse CSR path plus its popcount -- slower than not setting the hint at all, with no diagnostic. Enforce the documented range with RAFT_EXPECTS. NaN is rejected too, since it compares false against both bounds. Behavior for in-contract values, including the negative default, is unchanged. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Contributor
Author
|
@lowener this is number 2 in my series after 1) the dot product normalization constant optimization and before 3) the scatter gather and threshold fixing |
PrefilteredBruteForceInputs::sparsity is the density of kept entries (n_edges = m * n * sparsity), while search_params::filtering_rate is 1 - nnz / (n_queries * n_dataset). RunWithFilteringRateHint() passed the former as the latter, so the hint was the complement of the true sparsity and every hinted run took the dense path -- the lazy-nnz CSR branch was never exercised. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Contributor
Author
|
Pushed
With the sign fixed, the existing Also bumped the file's SPDX year to 2026; it was the only one of the five touched files still at 2025, so @lowener this PR hasn't had a CI run yet — could I get an |
Contributor
Author
|
@lowener merged This PR still hasn't had a CI run at all — could I get an |
rapids-bot Bot
pushed a commit
that referenced
this pull request
Aug 10, 2026
…spatch thresholds (#2321) ## The problem, in one picture `brute_force_search_filtered` chooses between a sparse (SDDMM) path and a dense post-filter path on one hardcoded test, `sparsity < 0.9f` — it runs SDDMM whenever **≤10% of the dataset passes the filter**. But SDDMM's cost climbs with the number of passing rows, while the dense path's is flat (it always does the full GEMM, then masks). The two curves cross long before 10%. cuVS rides the sparse path right past the crossover:  The pink track is what cuVS actually runs. At 10M×128 it is **12.6× slower than necessary** at the moment just before it finally gives up on SDDMM. The green curve is the path this PR adds. ## Two things are wrong, not one **1. The threshold is in the wrong place.** The true SDDMM/dense crossover is ~3%, not 10%. By the time the current rule hands off, SDDMM is already 2.2–4.4× slower than simply going dense. **2. The threshold is the wrong _kind_.** SDDMM competes against a fixed setup cost, not against work that scales with the dataset, so it stops paying off at an **absolute number of passing rows** — ~10k, roughly constant across the sweep. Expressed as a *fraction*, the correct handoff therefore **falls as the dataset grows**. At 10M rows it is ~0.1%, where the current rule waits until 10%: we keep SDDMM running up to 1,000,000 passing rows when it should hand off at ~10,000. **The error grows with the dataset.** ## Where each path actually wins  Blue is all that SDDMM should own. Everything left of the red line is what it owns today. The band between them — most of 0.1–10% — belongs to neither of the two paths cuVS currently has. ## The change Three paths, chosen by `select_filtered_search_path()`: - **`sddmm`** — unchanged; entered only below ~10k passing rows. - **`gather`** (new) — enumerate the passing rows, compact them (and their precomputed norms) into a dense `[n_pass x dim]` matrix, run the ordinary **unfiltered** search over it, map indices back. Cost scales with `n_pass` instead of `n_dataset`. - **`dense`** — unchanged; entered above the gather/dense crossover. Built from existing primitives — `thrust::copy_if` over `bitset_view::test`, `raft::matrix::gather`, `brute_force_knn_impl`, `thrust::gather`. No new kernels. The gather path applies to **bitset** filters only, which pass the same rows for every query. A **bitmap** passes a different set per query, so there is no single set of rows to compact; bitmaps keep a two-way dispatch, but with the corrected ~3% threshold instead of the flat 10%. ## What it costs today  Penalty against an oracle that always picks the fastest of the three paths: | dispatch rule | mean | p90 | worst | |---|---|---|---| | `sparsity < 0.9` (today) | 1.65× | 5.29× | **12.56×** | | this PR | **1.01×** | 1.00× | 1.49× | ## Tests `NEIGHBORS_TEST --gtest_filter='*Prefiltered*'`: **180/180 pass.** Six params were added to `selectk_inputs` at `n_dataset=300000, dim=128`, sized to land on each dispatch path. Since these tests check results against a CPU oracle and would pass just as well if the new path never ran, I instrumented the dispatch and confirmed each one is actually reached: | test | selectivity | passing rows | path taken | |---|---|---|---| | bitset Result/0 | 0.10% | 300 | `SDDMM` | | bitset Result/1–4 | 7.15% | 21,439 | **`GATHER`** | | bitset Result/5 | 21.0% | 63,076 | `DENSE` | | bitmap Result/0 | 0.10% | — | `SDDMM` | | bitmap Result/1–5 | 7.0–20% | — | `DENSE` | The four `GATHER` cases cover L2Expanded, InnerProduct, L2SqrtExpanded and CosineExpanded — the last exercising the gathered-norms path — and all agree with the CPU oracle. Bitmaps correctly carry no passing-row count and never gather. ## On the constants They are empirical, and I want to be plain that they are not derived from first principles. I tried to build a cost model (dense GEMM + gather traffic + select_k) so the `dim` and `n_dataset` dependence would fall out of hardware properties; solving it against the measured crossovers produces a **negative** gather-traffic coefficient, i.e. the model is wrong. So these are fits and I have not dressed them up as anything else. Instead I kept them **round and checked that they tolerate being wrong**. Perturbing any constant by 2× — roughly the spread I would expect across GPUs — leaves the mean penalty at 1.01–1.05× and the worst case at or below 2.7×. Still far better than the rule it replaces, which is 12.6× worst case *on the hardware it was presumably tuned for*. An earlier revision of this branch used 3-significant-figure power-law fits (`dim^0.4`, `dim^0.6`, …); they scored **identically** (1.01× mean), so the precision bought nothing and I removed it. If maintainers would rather these be tunable via `search_params`, or auto-calibrated once per device, both are easy from here. **Method:** fp16, L2Expanded, k=64, 100 queries, shared bitset; selectivity 0.1–60% × `n_dataset` ∈ {200K, 500K, 1M, 3M, 10M} × `dim` ∈ {128, 256, 512, 1024}; 100 reps on an idle RTX 5090 (sm_120), medians (cross-validated against an independent sweep to within 1.6–2.5%). ## Known gaps 1. **The thresholds are calibrated for `dim >= 128` and `n_dataset >= 200K`.** cuVS runs well below both — this repo's own prefiltered tests use `dim` from 1 to 2052 at `n_dataset=8192`. The fits are clamped to the measured range rather than extrapolated, but that corner is unmeasured. All 180 tests pass there; I simply cannot claim the *choice* is optimal. 2. **One pre-existing test config changes path**: `n_dataset=8192, dim=8, 10% density` moves SDDMM → dense. Both paths are correct (the test passes), but whether that flip is a fix or a regression is exactly what (1) has no data for. 3. **Bitmaps were never benchmarked.** The ~3% sparse/dense threshold is carried over from bitset measurements. The cost structure is the same shape, but a bitmap's CSR touches different columns per query, so SDDMM's locality is worse and the true crossover is probably *lower* than 3% — the error is in the safe direction, but it is inferred, not measured. 4. Single GPU architecture. 5. Touches the same function as #2120 (`filtering_rate` → base `search_params`). No conflict against `main` today, but the two will need reconciling if both land. ## Possible follow-up A bitmap could get a gather path too, by compacting the **union** of all queries' passing rows and applying the per-query mask inside that smaller matrix. It degenerates to exactly this PR's path when the filter is a bitset. Whether it pays off depends on how much the queries' filters overlap, so it needs its own benchmark — out of scope here. Authors: - Max Buckley (https://github.com/maxwbuckley) - Corey J. Nolet (https://github.com/cjnolet) - Micka (https://github.com/lowener) Approvers: - Micka (https://github.com/lowener) URL: #2321
…ispatch NVIDIA#2321 replaced the single `sparsity < 0.9` dense/CSR split with `select_filtered_search_path`, which chooses between sddmm, gather and dense and needs two exact quantities: `nnz_h` to size the CSR, and `n_pass` (bitset row count) to size the compacted dataset. Resolution keeps this PR's rule that a supplied `filtering_rate` must not cost a popcount, and adapts it to three paths instead of two: - The hint now feeds `selectivity` (= 1 - filtering_rate) rather than the old `sparsity`, and for a bitset also supplies an estimated `n_pass` so path selection has a row count to work with. - Only the dense path can run on the hint alone. If the chosen path is sddmm or gather, popcount lazily and then re-run the decision with the exact numbers, so a bad hint can never route into gather with an `n_pass` that under-runs `k`. The dense path — the common case, and the one the hint was added for — still skips the kernel and the host sync. - Drops this PR's old lazy-popcount block inside the CSR branch, now subsumed by the single lazy count at the dispatch site. The hint tests inherit NVIDIA#2321's new `selectk_inputs` entries, so `ResultWithFilteringRateHint` now covers all three paths rather than just dense and CSR. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
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
3 participants
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.
Summary
filtering_ratepreviously lived oncagra::search_paramsonly. Moving it to the basecuvs::neighbors::search_paramsmakes it inheritable by every algorithm'ssearch_paramsand removes the duplicate-per-algorithm pattern that would follow if other prefiltered paths wanted the same knob.brute_force::search_paramsthroughdetail::searchandbrute_force_search_filtered. Whenparams.filtering_rate >= 0, the hint is used directly assparsityand the per-searchbitset_view::count(res)popcount + host sync is skipped on the dense path. The sparse CSR path still needs an exact non-zero count to size the matrix, so popcount runs lazily there.ivf_flat,ivf_pq,hnsw) ignore it; theirsearch_paramsinherit the field but nothing reads it.Closes #1960.
cc @cjnolet — this is the cleanup for the popcount-per-search overhead from #1960.
Why
Both
cagra::search(with defaultfiltering_rate=-1.0) andbrute_force::search(always) callbitset_view::count(res)on every search to derive selectivity. That's a GPU reduction kernel + host stream sync per call. Documented measured impact on RTX 5090 at 1M/d=128/50% pass rate is ~18% latency overhead for CAGRA; brute_force had no escape hatch at all because itssearch_paramswas an empty derivation of the base.brute_force::search_paramswas already empty (struct search_params : cuvs::neighbors::search_params {};), so the field arrives via inheritance — no algorithm-specific addition needed.What changed
cpp/include/cuvs/neighbors/common.hppfloat filtering_rate = -1.0to basesearch_paramswith docstringcpp/include/cuvs/neighbors/cagra.hppcpp/src/neighbors/brute_force.cuparamsthrough the 4 macrosearch()entry pointscpp/src/neighbors/detail/knn_brute_force.cuhparamsarg todetail::searchandbrute_force_search_filtered; honor hint when>= 0, lazy popcount on CSR pathcpp/tests/neighbors/brute_force_prefiltered.cuResultWithFilteringRateHinttest cases for bitmap and bitset fixtures (float and half) reusing the existing parameter matrixTest plan
cpp/src/neighbors/brute_force.cucompiles cleanlycpp/tests/neighbors/brute_force_prefiltered.cucompiles cleanlyNEIGHBORS_TESTto confirmResultWithFilteringRateHintmatchesResultacross all parameter casesNotes
search_paramswassizeof == 1(empty); nowsizeof == 4(onefloat). All algorithm-specificsearch_paramsinherit, so their layout grows by 4 bytes. Appropriate for a release-cycle change.filtering_rateis currently only referenced from C++.filtering_rate < 0(default): same popcount, same downstream logic.🤖 Generated with Claude Code