perf(brute_force): add scatter-gather path and fix filtered-search dispatch thresholds

[maxwbuckley]

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 Promote filtering_rate to base search_params; honor it in brute_force #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.

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[maxwbuckley]

@lowener greatly appreciate your advice on this one, the third in a series after the 1) minor dot product optimizations and 2) the search params hint :)

[cjnolet]

/ok to test 48da83b

[maxwbuckley]

@cjnolet friendly ping — the /ok to test 48da83b above predates the two commits that landed after it (dd62cb6, SPDX headers; 0564963, merging main), so no build or test jobs have run against the current head. Could I get a /ok to test 0564963 when you get a chance?

[lowener]

LGTM I will approve after comments are addressed.
Thanks for this contribution

[lowener]

Why return sddmm here? Gather could still be better according to the figure in your PR

[maxwbuckley]

Good catch — you're right, and by more than I'd assumed. I measured real (library) gather vs SDDMM head-to-head in the sub-10k-passing-rows corner (RTX 5090, fp16, L2, k=64, N=1M, confirmed against N=3M):

dim	SDDMM wins below	penalty at n_pass≈9500 (SDDMM vs gather)
128	~1.7k rows	0.65 vs ≤0.41 ms → ~1.6×
1024	~3.0k rows	1.11 vs ≤0.54 ms → ~2.0×

(gather values are upper bounds — measured at n_pass≥12k, extended down by monotonicity — so the true crossover is even lower.)

The reason the original 10k looked right: it was calibrated against the standalone bench_bf_middle_path proxy, which rebuilds the index every rep and so overstates gather's cost by ~30–45% (~0.3 ms). The real brute_force_search_gathered reuses the gathered norms and skips that rebuild. Against the real path the SDDMM/gather crossover is ~2–3k rows, not 10k.

Fixed in a14e08e: kSddmmMaxPassingRows 10000 → 2000, so the 2k–10k band now routes to gather. Kept it flat/round (the crossover is only mildly dim-dependent, ~1.7k–3k across an 8× dim range) rather than fitting a dim curve.

One thing this raises that I did not chase here: the SDDMM/dense thresholds may have been calibrated against the same proxy, so they could lean the same way. Happy to re-derive those against the real paths as a follow-up if you'd like.

[lowener]

/ok to test 0564963