perf(brute_force): add scatter-gather path and fix filtered-search dispatch thresholds by maxwbuckley · Pull Request #2321 · NVIDIA/cuvs

maxwbuckley · 2026-07-14T19:06:01Z

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

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.
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.
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.
Single GPU architecture.
Touches the same function as Add a filtering_rate hint to brute_force::search_params #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.

…h dispatch thresholds Filtered brute-force search picks between a sparse (SDDMM) path and a dense post-filter path on `sparsity < 0.9`, i.e. it runs SDDMM whenever <=10% of the dataset passes the filter. Two things are wrong with that. First, the threshold is in the wrong place. SDDMM's cost grows with the number of passing rows while the dense path's is flat, so across a sweep of selectivity x n_dataset x dim (fp16, k=64, 100 queries, RTX 5090), SDDMM is already 2.2-4.4x slower than dense by the time it hands off at 10%. The real SDDMM/dense crossover is ~2.5-4%. Second, and more importantly, the threshold is the wrong *kind*. SDDMM stops paying off at an absolute number of passing rows -- roughly constant in n_dataset -- not at a fraction of the dataset, because it competes against a fixed setup cost rather than against work that scales with n_dataset. Expressed as a fraction, the correct handoff therefore falls as n_dataset grows: at n_dataset=10M it is ~0.11%, where the current rule waits until 10%. The error grows with the dataset. This adds a third path for the band the other two leave uncovered. For a bitset filter (which selects the same rows for every query) the passing rows are enumerated, compacted into a dense [n_pass x dim] matrix along with their precomputed norms, searched with the ordinary unfiltered kernel, and the resulting indices mapped back. Its cost scales with n_pass rather than n_dataset, so it wins over most of the 0.1-10% band that SDDMM currently owns. Measured against the best of the three paths, the current dispatch is up to 12.6x slower than it needs to be (n_dataset=10M, dim=128, 10% selectivity). A bitmap filter selects different rows per query, so there is no single set of rows to compact; bitmaps keep a two-way dispatch, but with the corrected SDDMM/dense threshold rather than the flat 10%. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

copy-pr-bot · 2026-07-14T19:06:05Z

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.

…lerant constants The previous constants were power-law fits to a single GPU (dim^0.4, dim^0.6, dim^0.22, 9.3, 212500, 11000). That is false precision: the exponents are not derivable from a cost model -- solving one against the measured crossovers yields a negative gather-traffic coefficient -- so they are curve-fits, and three significant figures of curve-fit do not transfer to other hardware. Replace them with four round constants. Over the measured sweep this is not a regression: mean penalty against a perfect oracle is 1.01x either way, and worst case improves slightly (1.52x -> 1.49x). All the extra precision bought nothing. The simpler rule also degrades gracefully. Perturbing any constant by 2x, which is roughly the spread expected across GPUs, keeps the mean penalty at 1.01-1.05x and the worst case at or below 2.7x -- still far better than the 1.65x mean and 12.6x worst case of the `sparsity < 0.9` rule it replaces. Also drops the commentary explaining the fits, which no longer exist. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

Regret figures against the previous rule describe a code review, not the code. They tell a future maintainer of this function nothing, and they go stale the moment anyone retunes. They live in the PR description instead. Also drops a reference to a benchmark README that is not part of this branch. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

A bitmap passes a different set of rows per query, so "the rows it passes" is not a thing. The dispatch was manufacturing one anyway -- nnz / n_queries -- so that it could reuse a single variable, then dividing by n_dataset to recover a selectivity it already had. Integer division made that lossy: a bitmap passing 5 pairs across 100 queries gave a row count of 0, hence a selectivity of exactly zero. In general it understated a bitmap's selectivity and pulled borderline cases onto the sparse path. Pass the selectivity (well-defined for both filter kinds, and the quantity the sparse/dense choice actually turns on) and make the row count a std::optional that only a bitset supplies -- which is also the only filter that can take the gather path that needs it. Dispatch is unchanged except at the sparse/dense boundary for bitmaps, where the rounding used to be wrong. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

maxwbuckley · 2026-07-15T05:57:23Z

@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 · 2026-07-16T19:23:30Z

/ok to test 48da83b

…ILIATES Fixes the verify-copyright pre-commit failure in CI. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

maxwbuckley · 2026-07-23T09:07:40Z

@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 · 2026-07-23T12:38:39Z

+    return sddmm_beats_dense ? filtered_search_path::sddmm : filtered_search_path::dense;
+  }
+  const int64_t n_pass = *passing;
+  if (n_pass < kSddmmMaxPassingRows && sddmm_beats_dense) { return filtered_search_path::sddmm; }


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

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 · 2026-07-23T13:16:02Z

/ok to test 0564963

Address review: replace the four rmm::device_uvector scratch buffers in brute_force_search_gathered with owning raft::make_device_{vector,matrix}. The two 2D buffers (gathered, compact_neighbors) become device matrices, so `gathered.view()` replaces the manual make_device_matrix_view wrapping and compact_neighbors.size() replaces the n_queries*k flatten arithmetic. Drop the now-unused stream handle and the redundant dispatch comment. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Measured real (library) gather vs SDDMM in the sub-10k passing-row corner on an RTX 5090 (fp16, L2, k=64, N=1M/3M). Gather overtakes SDDMM at ~1.7k-3k passing rows, not 10k -- the earlier 10k estimate came from a standalone gather proxy that rebuilds its index each rep and so overstated gather's cost by ~30-45%. Lower kSddmmMaxPassingRows so the 2k-10k band routes to gather. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

…ltered-dispatch-thresholds

maxwbuckley · 2026-07-26T14:11:39Z

Merged main in (ee4bd18) to clear the Recently Updated gate — the branch is now up to date with base. The earlier /ok to test 0564963 predates this merge, so could I get a fresh /ok to test ee4bd18 to run build + tests against the current head? Thanks!

lowener · 2026-07-27T09:35:28Z

/ok to test ee4bd18

lowener

LGTM

lowener · 2026-07-29T16:56:21Z

Spotted 2 flaky tests:

[CI] CLUSTER_TEST kmeans flaky #2100
- KmeansFitBatchedTests/KmeansFitBatchedTestF.Result/3
[CI] NEIGHBORS_ANN_NN_DESCENT_TEST flaky #2093
- AnnNNDescentTest/AnnNNDescentTestUI8_U32.AnnNNDescent/222
  Re-running tests

lowener · 2026-07-30T09:27:50Z

/ok to test 0cf4d77

maxwbuckley · 2026-08-06T19:02:51Z

@lowener the branch had drifted behind main again, which re-armed the Recently Updated gate, so I've merged main in (764fa9d). The branch is up to date with base and the diff is unchanged from the approved state — still just the two files (+177/-8).

The last full run (0cf4d77) was green across the board after your flaky-test re-run. Could I get an /ok to test 764fa9d so it can go in? Happy to keep re-merging, but each round trip re-opens the same gate, so if it's possible to merge shortly after the next green run that would save us both the loop. Thanks!

lowener · 2026-08-07T16:07:21Z

/ok to test 764fa9d

lowener · 2026-08-10T09:34:33Z

/merge

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

… the SDDMM path (#2505) ## The problem On the sparse (SDDMM) branch of `brute_force_search_filtered`, the filter is converted to CSR twice per search: ```cpp auto csr = raft::make_device_csr_matrix<DistanceT, IdxT>(res, n_queries, n_dataset, nnz_h); std::visit([&](const auto& actual_view) { actual_view.to_csr(res, csr); }, *filter_view); // (1) ... raft::sparse::linalg::masked_matmul(res, queries, dataset_view, actual_view, csr_view); // (2) ``` `masked_matmul` builds a *sparsity-preserving* CSR over the caller's structure view and runs `bitset_to_csr` on it, which writes through to the same `indptr`/`indices` cuVS just filled: ```cpp auto C_matrix = raft::make_device_csr_matrix<output_t, index_t>(handle, compressed_C_view); raft::sparse::convert::bitset_to_csr(handle, mask, C_matrix); ``` So (2) recomputes (1), identically. ## The change Drop cuVS's `to_csr` call and let `masked_matmul` do the conversion once. ## Both conversions really were running `nsys` over 21 SDDMM-path searches (1M x 128, 0.5% selectivity), launch counts before → after: | kernel | before | after | |---|---|---| | `calc_nnz_by_rows_kernel` | 42 | 21 | | `fill_indices_by_rows_kernel<check_nnz=false>` | 21 | **0** | | `fill_indices_by_rows_kernel<check_nnz=true>` | 21 | 21 | | `repeat_csr_kernel` | 42 | 21 | | exclusive-scan (`DeviceScanKernel`) | 42 | 21 | | CSR value fill | 42 | 21 | | `cusparse::sddmm_ker` | 21 | 21 | | `csr_to_coo_kernel` | 21 | 21 | Every conversion kernel halves while `sddmm_ker` and `csr_to_coo` are unchanged, i.e. no actual work is removed. `check_nnz=false` is the sparsity-*owning* instantiation, which only cuVS's `to_csr` could produce; it disappears entirely. A host sync goes with it: `to_csr` on a sparsity-owning matrix reads `nnz` back to the host to size the matrix, which the sparsity-preserving path inside `masked_matmul` never does. 8-byte device-to-host copies over 29 searches drop from 50 to 29. ## Why relying on `masked_matmul` is safe Same problem run twice, once with C's structure pre-filled correctly and once with a valid but deliberately wrong pattern (same row counts, every column index 0): ``` m=4 nnz=112 path=faster_dot_on_csr indptr MATCH indices MATCH values MATCH m=16 nnz=448 path=sddmm indptr MATCH indices MATCH values MATCH ``` Bit-identical on both dispatch branches, so `masked_matmul` unconditionally derives C's sparsity from the mask. This is deliberate in RAFT — `bitset_to_csr` carries a `constexpr bool check_nnz = is_device_csr_sparsity_preserving_v<csr_matrix_t>;` branch for exactly this case — but it is not **documented**: `masked_matmul` marks `C` `[inout]` only with respect to its values, and the sparsity-preserving CSR it constructs is documented as "allowing the values in the sparsity to change but not the sparsity itself". Happy to hold this PR if reviewers would rather a RAFT doc fix landed first. *(An earlier revision of this description said RAFT's `masked_matmul` test depends on the overwrite. It does not — the test builds C's structure from the same mask it passes in, so it cannot discriminate. Hence the experiment above.)* ## Measurements Paired A/B, arms alternating within every round and every (N, dim) block so drift hits both equally: 10 rounds x 200 reps per config, separately built `libcuvs.so` swapped via `LD_LIBRARY_PATH`, dedicated idle RTX 5090 (sm_120, CUDA 12.9, WSL2), fp16 / L2Expanded / k=64 / 100 queries / shared bitset. CIs are Student-t over the 10 paired round-medians. | N | dim | selectivity | baseline | this PR | speedup | saving (95% CI) | t | |---|---|---|---|---|---|---|---| | 1M | 128 | 0.1% | 0.383 ms | 0.238 ms | **1.61x** | +136 us [93, 180] | 7.1 | | 1M | 128 | 0.5% | 0.548 ms | 0.407 ms | 1.35x | +138 us [106, 170] | 10.9 | | 1M | 128 | 1.0% | 0.708 ms | 0.579 ms | 1.22x | +133 us [115, 151] | 12.7 | | 1M | 128 | 3.0% | 1.409 ms | 1.272 ms | 1.11x | +139 us [123, 155] | 18.4 | | 1M | 1024 | 0.1% | 0.468 ms | 0.323 ms | **1.45x** | +146 us [109, 183] | 9.3 | | 1M | 1024 | 0.5% | 0.783 ms | 0.663 ms | 1.18x | +129 us [105, 153] | 12.3 | | 1M | 1024 | 1.0% | 1.239 ms | 1.082 ms | 1.15x | +147 us [125, 168] | 14.8 | | 1M | 1024 | 3.0% | 4.643 ms | 4.467 ms | 1.04x | +169 us [124, 214] | 6.7 | | 10M | 128 | 0.1% | 0.750 ms | 0.589 ms | 1.27x | +158 us [134, 182] | 15.9 | | 10M | 128 | 0.5% | 2.491 ms | 2.338 ms | 1.07x | +152 us [139, 165] | 26.5 | | 10M | 128 | 1.0% | 4.707 ms | 4.531 ms | 1.04x | +201 us [140, 261] | 8.2 | | 10M | 128 | 3.0% | 14.084 ms | 13.760 ms | 1.02x | +335 us [252, 419] | 7.9 | All 12 SDDMM-path configs improve with the interval excluding zero. The saving is a roughly **fixed ~150 us per search** rather than a ratio, so it is largest where the search is cheapest — which post-#2321 is exactly where the sparse path lives, below ~2000 passing rows. **Null control.** Dense-path configs cannot be affected and are flat: +1.8 us [-11.3, +15.0] at 1M/128/25%, -57.6 us [-134.0, +18.9] at 1M/1024/25%, +96.7 us [-277.8, +471.1] at 10M/128/25%. All three span zero. 10M x 1024 is drift-dominated on this host (±2 ms swings in both arms, including the control) and is excluded rather than reported. **The saving is not the removed sync.** Removing the *other* per-search sync (`bitset_view::count`, fed instead from an environment variable, results bit-identical) is worth median +0.0 us, 0 of 15 configs significant — a sync on an already-idle stream costs nothing. The win is the removed conversion work plus the pipeline bubble from draining the stream immediately before a block of host-side cuSPARSE setup. ## Correctness - Top-k ids and distances **bit-identical** to the unpatched build across 14 configs (dim ∈ {128, 1024} x selectivity ∈ {0.1%, 0.5%, 1%, 3%, 5%, 10%, 25%}), spanning both dispatch paths. - `NEIGHBORS_TEST --gtest_filter='*Prefiltered*'`: **156/156 pass**. That run is on a pre-#2321 checkout, so 39 parameterisations against 45 on `main`; the 6 that #2321 added cover the gather and dense paths, which this change does not touch. - Compiles clean against `main` with the build's own `-Werror=all-warnings`. ## RAFT-side follow-ups, not in this PR - `masked_matmul` allocates `nnz * sizeof(output_t)` for `C_matrix`'s element buffer that nothing ever reads — both `sddmm` and `faster_dot_on_csr` write through `C.get_elements()`. Memory-only, no measurable latency change. - `masked_matmul` should document that it overwrites `C`'s sparsity pattern from the mask, and its test should pass a structure that differs from the expected one so it actually covers the behaviour relied on here. Authors: - Max Buckley (https://github.com/maxwbuckley) - Micka (https://github.com/lowener) Approvers: - Micka (https://github.com/lowener) URL: #2505

github-project-automation Bot added this to Unstructured Data Processing Jul 14, 2026

maxwbuckley and others added 3 commits July 14, 2026 21:16

maxwbuckley mentioned this pull request Jul 14, 2026

[FEA] Roaring Bitmap support #1972

Closed

maxwbuckley marked this pull request as ready for review July 14, 2026 21:28

maxwbuckley requested a review from a team as a code owner July 14, 2026 21:28

cjnolet assigned maxwbuckley Jul 16, 2026

cjnolet added improvement Improves an existing functionality non-breaking Introduces a non-breaking change labels Jul 16, 2026

cjnolet added this to Unstructured Data Processing Jul 16, 2026

cjnolet moved this to In Progress in Unstructured Data Processing Jul 16, 2026

maxwbuckley and others added 2 commits July 16, 2026 22:16

style: update SPDX copyright headers to 2026 NVIDIA CORPORATION & AFF…

dd62cb6

…ILIATES Fixes the verify-copyright pre-commit failure in CI. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Merge branch 'main' into perf/brute-force-filtered-dispatch-thresholds

0564963

lowener reviewed Jul 23, 2026

View reviewed changes

maxwbuckley and others added 3 commits July 24, 2026 12:35

Merge remote-tracking branch 'upstream/main' into perf/brute-force-fi…

ee4bd18

…ltered-dispatch-thresholds

lowener approved these changes Jul 27, 2026

View reviewed changes

Merge branch 'main' into perf/brute-force-filtered-dispatch-thresholds

0cf4d77

Merge remote-tracking branch 'upstream/main' into HEAD

764fa9d

maxwbuckley mentioned this pull request Aug 6, 2026

Add a filtering_rate hint to brute_force::search_params #2120

Open

rapids-bot Bot merged commit abfb8b9 into NVIDIA:main Aug 10, 2026
162 of 166 checks passed

github-project-automation Bot moved this to Done in Unstructured Data Processing Aug 10, 2026

maxwbuckley mentioned this pull request Aug 25, 2026

perf(brute_force): drop the duplicate bitset/bitmap→CSR conversion on the SDDMM path #2505

Merged

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×

Conversation

maxwbuckley commented Jul 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

The problem, in one picture

Two things are wrong, not one

Where each path actually wins

The change

What it costs today

Tests

On the constants

Known gaps

Possible follow-up

Uh oh!

copy-pr-bot Bot commented Jul 14, 2026

Uh oh!

maxwbuckley commented Jul 15, 2026

Uh oh!

cjnolet commented Jul 16, 2026

Uh oh!

maxwbuckley commented Jul 23, 2026

Uh oh!

lowener left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

lowener Jul 23, 2026

Choose a reason for hiding this comment

Uh oh!

maxwbuckley Jul 24, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

lowener commented Jul 23, 2026

Uh oh!

maxwbuckley commented Jul 26, 2026

Uh oh!

lowener commented Jul 27, 2026

Uh oh!

lowener left a comment

Choose a reason for hiding this comment

Uh oh!

lowener commented Jul 29, 2026

Uh oh!

lowener commented Jul 30, 2026

Uh oh!

maxwbuckley commented Aug 6, 2026

Uh oh!

lowener commented Aug 7, 2026

Uh oh!

lowener commented Aug 10, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

maxwbuckley commented Jul 14, 2026 •

edited

Loading