[CUB] Strengthen DeviceHistogram benchmarks and correctness coverage by robobryce · Pull Request #10548 · NVIDIA/cccl

robobryce · 2026-07-30T10:49:12Z

Why

The DeviceHistogram optimization work in #10547 needs a trustworthy, reusable way to generate difficult inputs, detect incorrect output, and compare algorithm behavior before implementation changes can be reviewed independently.

The previous benchmark mostly measured a narrow family of inputs, and some invalid benchmark cells could be skipped rather than fail. That made it too easy for an optimization to look successful because it recognized the benchmark pattern, silently avoided an unsupported case, or produced an incorrect histogram.

This draft is the first focused extraction from the raw autoresearch branch. It contains benchmark and correctness infrastructure only; it does not include the new shared-memory or high-bin caching algorithms. The next focused PR will be stacked on this branch while it is under review, then retargeted to main after this lands.

What changed

Reproducible input shapes

The histogram benchmarks now generate explicit, deterministic input shapes instead of relying on a single entropy-style knob. The available shapes exercise materially different workloads, including concentrated hot bins, uniform traffic, and distributions intended to stress cache capacity and collisions. Single-channel and multi-channel benchmarks use the same input contracts.

The benchmark documentation and Python tooling describe and visualize those distributions, run comparable sweeps, and retain the configuration needed to interpret a result. The plotting paths use current algorithm keys so results from later implementation PRs can be compared without rewriting the analysis scripts.

Independent correctness checks

Benchmark results are checked bin by bin against a reference implementation. A mismatch is now fatal instead of silently skipping the cell. RANGE benchmarks use genuinely non-uniform integer levels, and the reference path is deliberately separate from the device bin-computation path so the same defect is less likely to affect both sides.

The tests cover the input-shape generators and the thread-local benchmark cache. Existing DeviceHistogram coverage is extended for the boundary cases exposed while building the harness.

Wide and overflow-prone cases

The benchmark can be built with alternate counter and offset types, including a 64-bit baseline. It covers larger bin counts, including 16,384 and 60,000 bins, and uses the full representable integer sample range where applicable.

This work also fixes correctness and undefined-behavior issues found by the new coverage:

signed overflow in ScaleTransform::ComputeBin for negative minimum levels;
bin computation when the level range is wider than LevelT;
row-stride overflow in invalid multi-channel benchmark cells;
stale persisting-L2 state affecting comparisons between benchmark iterations.

Scope and follow-up

This is intentionally a benchmark/testing foundation rather than an algorithm PR. It is larger than a normal test-only change because the generators, reference validation, C++ benchmark integration, and analysis tools form one reproducibility contract.

The next focused PR will add the shared-memory privatized enhancements as a stacked change. Its intended scope is one generic dynamic-shared-memory privatized kernel, byte- and counter-width-aware capacity checks, occupancy-aware launch handling, conservative selection, and focused correctness tests. It will exclude the new high-bin cache, direct-atomic selector experiments, and the historical staging/combine designs from #10547.

Later PRs will separately cover the high-bin global-memory baseline, the cache algorithm, selector/tuning policy, and C Parallel integration, as outlined in #10547.

Validation

Formatting and repository checks:

pre-commit run --files <all changed files>
git diff --check upstream/main...HEAD

Both pass.

Configured and built with CUDA 13.3.33 for SM90:

CUDAHOSTCXX=/usr/bin/c++ CXX=/usr/bin/c++ \
  cmake --fresh --preset cub-cpp20 -DCMAKE_CUDA_ARCHITECTURES=90

ninja -C build/cub-cpp20 \
  cub.test.device.histogram_input_shapes \
  cub.test.device.histogram_thread_local_cache \
  cub.test.device.histogram

All requested targets compile successfully.

Runtime execution is currently blocked by the available test environment. The machine has an NVIDIA B200, but its installed driver rejects PTX produced by CUDA 13.3.33 with cudaErrorUnsupportedPtxVersion: the provided PTX was compiled with an unsupported toolchain. The same startup failure affects the existing histogram test and both new test binaries, before any test assertion runs. Runtime results will need to be collected with a compatible driver/toolkit pair.

Relationship to the research snapshot

This code was extracted and formatted from the raw autoresearch output in #10547. Scratch-directory changes, implementation experiments, generated artifacts, design notes, and the new histogram algorithms are not included here.

Three changes to cub/benchmarks/bench/histogram/{even,range}.cu so the benchmarks exercise the code paths real users hit: - range.cu: build levels[] with quadratic spacing (still strictly monotonic across [lower_level, upper_level]) so DispatchRange stays on the SearchTransform path. The previous thrust::sequence boundaries were perfectly uniform, letting any uniform-detection fast path collapse the bench to DispatchEven performance. - both: replace the power-of-two Elements{io} axis with non-power-of-two sizes so tunings that hard-code on round counts (exact tile multiples, pow2 shortcuts) get measured at sizes where those shortcuts cannot fire. The total axis cardinality is unchanged. - both: switch to the manual-timer exec_tag and call cudaCtxResetPersistingL2Cache() outside the timed window. nvbench's cold measurement already evicts cached lines between iterations, but it does not demote persistence-marked addresses set via cudaStreamSetAttribute / cudaAccessPolicyWindow.

…cache coverage Benchmarks - multi/even.cu, multi/range.cu: same hardening as the prior commit's even.cu/range.cu changes — quadratic-spaced range levels (still strictly monotonic across [lower_level, upper_level]) so DispatchRange stays on the SearchTransform path; the manual-timer exec_tag with cudaCtxResetPersistingL2Cache() outside the timed window; non-power- of-two Elements{io} so tunings that hard-code on round counts (exact tile multiples, pow2 shortcuts) are exercised at sizes where those shortcuts cannot fire. Axis cardinality is unchanged. - even.cu, range.cu, multi/{even,range}.cu: replace two of the four Bins values (128 -> 100, 2048 -> 2000) so tunings that hard-code on power- of-two bin counts cannot use those shortcuts. Cardinality unchanged. Tests - catch2_test_device_histogram.cu: setup_bin_levels_for_range now perturbs interior levels by +/- min_bin_width/4 (alternating sign), falling back to uniform when the type is too tight (e.g. byte-sample with 256 levels). The std::upper_bound reference already handled arbitrary spacings; this just exercises the SearchTransform code path in addition to the uniform-detection fast path. - catch2_test_device_histogram_thread_local_cache.cu (new): three Catch2 cases targeting the thread_local detection_stream / detection_ buf cache in dispatch_range. Sequential calls across multiple user streams, four-thread concurrent calls on the same device, and a single-thread cross-device case that skips when fewer than two GPUs are present.

…eration The default cudaLimitPersistingL2CacheSize is 0, so hardcoding 0 (rather than relying on the default) defends against prior benchmarks in the same nvbench process having bumped the reservation. The cudaCtxResetPersistingL2Cache call already demoted persistence-marked addresses; this extends the defense to the reservation itself.

Quadratic spacing produced bin widths spanning ~2n× (last bin vs first), which is not representative of typical workloads. Jittered uniform spacing (±25% of step, fixed mt19937 seed) keeps consecutive widths within ~3× while still defeating uniform-spacing detection so DispatchRange stays on the SearchTransform path.

Extends the Bins axis to cover the 10k-65k range, which sits between the existing 2000 and 2097152 entries. Applied to range/even and the multi-channel variants so all four histogram benches share the same axis.

Each cell of the four `cub.bench.histogram.{even,range,multi.even, multi.range}.base` benchmarks now runs the dispatch once before NVBench's timed window and compares the produced per-channel histogram bin-by-bin against an independent reference computed on-device with `thrust::for_each` + global `atomicAdd`. The warmup also checks the dispatch return code so a non-`cudaSuccess` return is reported instead of being silently discarded. The verifier runs entirely outside `state.exec`, so timed-region bandwidth is unchanged within measurement noise. Wall-clock per benchmark cell increases proportionally to the input size of that cell (the reference loops over every sample once on device). The verifier is on by default and can be disabled at run time by setting the environment variable `CUB_BENCH_HISTOGRAM_VERIFY` to one of: `0`, `false`, `no`, `off` (case-insensitive). Disabling it skips the warmup dispatch, the reference build, and the bin-by-bin compare. The verifier catches two bug classes that the existing CTest histogram suite does not: - dispatch-time errors (e.g. `cudaErrorInvalidValue` from a temp- storage size mismatch in the chunked-staging path) that are not reported by NVBench because the dispatch return code is dropped on the floor. - per-bin count corruption that still produces a non-empty histogram with the right shape but the wrong values (e.g. a partition mask that drops samples that should have landed in another partition's write set). These pass any sum-of-counts sanity check.

The jittered-uniform level construction in the range benches sets upper_level via get_upper_level, which previously returned num_bins for integer SampleT. That produced step = 1.0, so the ±0.25*step jitter sat in [-0.25, 0.25] and was annihilated by the integer cast in the level loop. The subsequent dedup-by-1 step then forced every collision back onto the next consecutive integer, leaving the level array bit-identical to a perfect uniform stride-1 sequence. A DispatchRange uniform-spacing detection then has nothing to detect against: it sees a perfectly uniform level array on every integer axis row and routes straight to the EVEN classify path - exactly the fast path the range bench is supposed to avoid measuring. Widen upper_level to ~4 * num_bins for integer SampleT so step is at least ~4 and ±step/4 jitter survives integer truncation as ±1, which is enough to break uniformity. Clamp to the type max when 4 * bins overflows SampleT; those axes (e.g. int8_t with bins >= 64) already have step < 1 and the level array is degenerate regardless of jitter.

…velT Two latent bugs in cub::DeviceHistogram surfaced when widening the bench to use the full SampleT range (lower_level = numeric_limits<SampleT>::min() for signed integers). 1. ScaleTransform stored `m_max`, `m_min`, `m_scale.fraction.range`, and `m_scale.fraction.bins` in `CommonT = common_type<LevelT, SampleT>`, then ComputeBin promoted through the wider `IntArithmeticT` only at the multiply/divide step. For narrow integer CommonT (int8_t, int16_t) the precomputed `range = max - min` overflowed CommonT before the promotion: int8_t with [-128, 127] gave `range = 255` truncated back to int8_t = -1, sign-extended in IntArithmeticT to 0xFFFFFFFF, and ComputeBin's division by that gigantic divisor returned 0 for every sample. The histogram was non-empty but every count landed in bin 0. Fix: introduce FractionStorageT = IntArithmeticT for integer CommonT (CommonT for non-integer types) and store both `range` and `bins` in it. Compute `max - min` through ULevelT = make_unsigned_t<T>: the intermediate cast is required because C++ integer promotion lifts `(uint8_t) - (uint8_t)` to int(127 - 128) = -1, and going directly to FractionStorageT sign-extends that to a huge garbage value. Truncating through ULevelT first lets unsigned modular wrap-around recover the correct difference. 2. The MayOverflow precondition check at the byte-sample EVEN dispatch sites in DispatchEven cast `num_levels - 1` to CommonT before passing it to MayOverflow: `static_cast<int8_t>(128) = -128` for int8_t, sign-extended in IntArithmeticT to 0xFFFFFF80, and the subsequent division `numeric_limits<IntArithmeticT>::max() / 0xFFFFFF80 = 1` reported overflow for any non-trivial range. Fix: pass `num_levels - 1` directly (it's already an `int`) and apply the same unsigned-promotion- safe subtraction in MayOverflow's `(upper - lower)` computation. 3. PassThruTransform::BinSelect computed `bin = static_cast<int>(sample)` for the byte-sample privatized histogram. For signed int8_t samples this preserved the sign, producing negative bin indices in [-128, -1] for half the input range; the kernel's `if (bin >= 0)` check then silently dropped them. Fix: cast through make_unsigned_t<_SampleT> first so int8_t(-128..127) reinterprets as uint8_t(128..255, 0..127). The existing "DeviceHistogram::HistogramEven num_bins exceeds LevelT range" test was asserting `cudaErrorInvalidValue` for inputs that are now correctly handled. Updated to assert success — the bin width can be fractional (smaller than one distinct LevelT value), and the integer ComputeBin path handles that without overflow once the storage-type and cast bugs are fixed. Co-Authored-By: Claude Opus 4.7 (1M context) <[email protected]>

For signed integer SampleT the bench now picks `lower = numeric_limits<SampleT>::min()` instead of `0`. This doubles the testable range — int8_t goes from 128 distinct values [0, 127] to 256 distinct [-128, 127], int16_t from 32768 to 65536 — letting the benchmarks exercise more bin counts before hitting the distinct-level-values cap. With this change and the matching DeviceHistogram fix, int8_t now runs the dense matrix at bins=128 and bins=255 (previously skipped or producing zero-filled histograms). Helpers added in histogram_common.cuh: - get_lower_level<SampleT>() returns numeric_limits::min() for signed integer SampleT and 0 otherwise. - max_representable_bins<SampleT>() returns the count of distinct SampleT values minus 1 (the upper bound on bins + 1 strictly-monotonic levels). For 64-bit and floating-point SampleT it's int64_t::max(), effectively unbounded for the bench's bin axes. The four bench files (`even`, `range`, `multi/even`, `multi/range`) swap their hardcoded `lower_level = 0` for `get_lower_level<SampleT>()` and gate on `num_bins > max_representable_bins<SampleT>()`. Co-Authored-By: Claude Opus 4.7 (1M context) <[email protected]>

`MultiHistogramEven`/`MultiHistogramRange` internally compute `row_stride_samples = elements * num_channels` and pass it as `OffsetT`. For `OffsetT = int32_t` and `num_channels = 4`, this caps usable elements at `INT_MAX / 4` (~536M); above that the cast wraps to a negative value and the kernel produces zero output without raising an error. The bench correctness check catches the empty histogram, but the skip reason ("opt=0 ref=N") obscures the underlying overflow. Add an explicit overflow check in the multi-channel benches so cells that would hit this limit skip cleanly with a descriptive reason. This matters at autocuda matrix axes >= 1G elements: with three active channels the row stride becomes `3 * 1G = 3G`, well above `INT_MAX`. The single-channel benches don't need this check; their `row_stride = elements` and elements is already bounded by the `int64_t` axis type. Adding I64 OffsetT to the multi-channel type list (a separate change) would lift this restriction. Co-Authored-By: Claude Opus 4.7 (1M context) <[email protected]> (cherry picked from commit 99fa749)

…silent skip The per-cell bin-by-bin verifier previously signaled a mismatch by throwing std::runtime_error from the benchmark body. nvbench catches that and marks the cell `Skipped: Yes`, then exits 0 -- so a kernel that computes wrong per-bin counts on the hard cells had those cells silently dropped from the geomean, which INFLATED the reported bandwidth (a reward-hacking hole). Replace the verifier throws with bench_fatal(), which prints the diagnostic and std::abort()s so the binary exits non-zero and the whole trial fails loudly. Legitimate skips (row-stride overflow) use state.skip(...) and are unaffected. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]> (cherry picked from commit 8d821a7)

The histogram benchmarks built their input via the shared nvbench_helper generate(elements, entropy, lower, upper), whose bitwise-AND "entropy" knob is non-linear (bunched at the extremes), always pins the hot bin to the zero value, and cannot express multi-hot or cache-adversarial inputs. Add cub/benchmarks/bench/histogram/histogram_inputs.cuh: shapes are decided in bin-index space then mapped to sample values (EVEN: bin midpoint; RANGE: level-interval midpoint), so the existing in-bench verifier validates every shape automatically. The Entropy string axis is replaced by an InputShape axis whose values carry an optional inline knob "name:value": * concentrated:E -- spike-slab family, E = target normalized entropy. E=1.0 is exact uniform (equal-count tiling, zero count variance), E=0.0 is constant fill, in between is one scattered hot bin over a uniform floor. Replaces and generalizes the old entropy sweep, continuously. * powerlaw:E -- decaying warm set; rank exponent solved for target entropy E (an independent knob from concentrated). * zipf:s -- decaying warm set, classic exponent s. * hash_synonym:h -- hot bins collide on one cache slot (attacks hashed cache). * capacity_cliff:m-- m * cache_slots equiprobable bins (attacks bounded cache). * stale_resident:m-- cold prefix claims slots, then a hot bulk (attacks no-evict). * temporal_phases:n, strided_sweep:n -- ordering-structured adversaries. The hot bin is scattered off zero via a fixed coprime permutation, so the mode is no longer always bin 0. Add catch2_test_device_histogram_input_shapes.cu validating each shape's bin distribution / ordering and the monotonicity of the entropy knobs (200k+ assertions). All four bench binaries build and run clean across every shape with the in-bench verifier on (no correctness aborts). Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]> (cherry picked from commit 79dfdc4e949c939d5ec97679cbfa1d184bdd50e7)

…min_level The integral ComputeBin path computed `sample - min_level` in the signed sample type T before casting to the wider unsigned IntArithmeticT. For signed T with a sufficiently negative min_level (e.g. T=int32_t, min_level=INT_MIN), the signed subtraction overflows (undefined behaviour); on two's complement it wraps negative and the subsequent widening produces a wrong magnitude, so the kernel computes a garbage bin index and the sample is dropped from the output histogram. Top-of-range samples in particular were lost, producing small undercounts at bin = num_bins - 1. Fix: compute the difference via the unsigned representation of T (mirroring ScaleTransform::ComputeScale's `max_level - min_level`), which wraps modularly to the correct non-negative difference, then widen to IntArithmeticT. Backport of 0884164 onto main. The original sat atop later EVEN-path optimizations (a magic-multiplier `range_divider` and a `bins_eq_range` fast path) that are not present on main; this commit applies ONLY the overflow fix, keeping main's `* bins / range` integer division unchanged. (cherry picked from commit 0884164) Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>

Benchmark-only tooling backport (no library/dispatch changes): - Input-shape generator rework: histogram_inputs.cuh gains the sawtooth shape and the random-order uniform endpoint, redefines concentrated (random/entropy) and stale_resident (cache-thrash), and drops capacity_cliff; even/range/multi axis lists updated to match. Bench-only -- no dispatch/kernel code is touched. - histogram_input_design.py: bit-exact host mirror of the generators (shared module). - histogram_input_characterization.py: per-shape characterization figures (distribution / rank-frequency / position-in-sequence). - histogram_algo_perf.py: per-shape GiB/s-vs-#bins figures with a log-y axis and the selector-default + (optional) upstream-baseline reference series. - histogram_algo_sweep.py: reproducible perf-sweep driver. (Algorithm forcing via CUB_HISTO_FORCE_ALGO is a no-op on stock dispatch -- the forced columns collapse onto `default` here; the hook lives with the experimental optimization work.) - README_plots.md: documents the scripts and the sweep/plot workflow. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>

Sync histogram_algo_perf.py so the cache-hit-rate panels read the current direct_cuckoo / direct_single_probe keys (the earlier tooling backport carried the pre-rework direct_atomic_* spellings). No behavior change without hit-rate data, but keeps the plotter consistent with histogram_hitrate_sweep.py. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>

… baseline build Add TUNE_CounterT / TUNE_OffsetT guards (inert when undefined -> baseline dispatch unchanged) so a .base.u64 variant of this baseline can be built, giving the feature branch's unified 64-bit-counter sweep a fair as-shipped `main` baseline series. Co-Authored-By: Claude Opus 4.8 (1M context) <[email protected]>

copy-pr-bot · 2026-07-30T10:49:16Z

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.

robobryce · 2026-07-30T12:54:03Z

Superseded by the same focused change under the corrected branch namespace: #10555.

robobryce · 2026-07-30T12:54:57Z

Replacement draft under the corrected branch namespace: #10555 ().

robobryce · 2026-07-30T12:55:13Z

Corrected branch name: pr/histocache/benchmark-testing.

brycelelbach and others added 17 commits July 30, 2026 10:05

[cub] Add 16384 and 60000 bin sizes to histogram benchmarks

1841b47

Extends the Bins axis to cover the 10k-65k range, which sits between the existing 2000 and 2097152 entries. Applied to range/even and the multi-channel variants so all four histogram benches share the same axis.

[cub] Format histogram benchmark and test additions

7f27347

github-project-automation Bot added this to CCCL Jul 30, 2026

github-project-automation Bot moved this to Todo in CCCL Jul 30, 2026

cccl-authenticator-app Bot moved this from Todo to In Progress in CCCL Jul 30, 2026

robobryce mentioned this pull request Jul 30, 2026

[CUB] Add runtime-sized shared-memory histogram privatization #10553

Closed

robobryce closed this Jul 30, 2026

robobryce deleted the pr/histogram-bench-testing branch July 30, 2026 12:54

robobryce mentioned this pull request Jul 30, 2026

[cub] Add static and runtime-sized shared-memory histogram privatization #10556

Open

Navigation Menu

Uh oh!

[CUB] Strengthen DeviceHistogram benchmarks and correctness coverage - #10548

[CUB] Strengthen DeviceHistogram benchmarks and correctness coverage#10548
robobryce wants to merge 17 commits into
NVIDIA:mainfrom
robobryce:pr/histogram-bench-testing

robobryce commented Jul 30, 2026

Uh oh!

copy-pr-bot Bot commented Jul 30, 2026

Uh oh!

robobryce commented Jul 30, 2026

Uh oh!

robobryce commented Jul 30, 2026

Uh oh!

robobryce commented Jul 30, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

robobryce commented Jul 30, 2026

Why

What changed

Reproducible input shapes

Independent correctness checks

Wide and overflow-prone cases

Scope and follow-up

Validation

Relationship to the research snapshot

Uh oh!

copy-pr-bot Bot commented Jul 30, 2026

Uh oh!

robobryce commented Jul 30, 2026

Uh oh!

robobryce commented Jul 30, 2026

Uh oh!

robobryce commented Jul 30, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants