[cub] Strengthen DeviceHistogram benchmarks and correctness coverage - #10555
Open
robobryce wants to merge 21 commits into
Open
[cub] Strengthen DeviceHistogram benchmarks and correctness coverage#10555robobryce wants to merge 21 commits into
robobryce wants to merge 21 commits into
Conversation
Contributor
Contributor
|
Split the benchmarking scripts into a new repo called robobryce/histocache-benchmarking-scripts; leave it out of this PR. |
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
brycelelbach
reviewed
Aug 8, 2026
5 tasks
robobryce
changed the title
Author
|
Thanks. No further action is needed on those two low-value refactor suggestions; the final branch and validation remain unchanged. |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
🧹 Nitpick comments (4)
cub/test/catch2_test_device_histogram_input_shapes.cu (1)
27-29: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: Declare
d_inputandh_inputasconst. This helper only copies or reads these vectors before it returns. As per coding guidelines, “All variables that are not modified must be declaredconst.”Source: Coding guidelines
cub/benchmarks/bench/histogram/even.cu (1)
7-7: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: Use canonical angle-bracket include paths at all sites. Quoted header inclusions violate the required inclusion form.
cub/benchmarks/bench/histogram/even.cu#L7-L7: Replace the quotedhistogram_inputs.cuhinclude.cub/test/catch2_test_device_histogram_input_shapes.cu#L15-L17: Replace the quoted project header includes.cub/benchmarks/bench/histogram/range.cu#L10-L11: Replace the quoted histogram header includes.As per coding guidelines, “All header inclusions must use angle-bracket syntax.”
Source: Coding guidelines
cub/test/catch2_test_device_histogram.cu (2)
532-532: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: Use the required
constexprandconstqualifiers.
max_level_countdepends only onsizeof(sample_t), so both declarations can useconstexpr auto.d_samplesis not modified in this scope, so useconst auto. Keepd_histo_outmutable because the histogram call writes through its buffer.As per coding guidelines, “All variables that are not modified must be declared
const” and “All variables that can be evaluated at compile time must be declaredconstexpr.Also applies to: 542-542, 774-774
Source: Coding guidelines
775-780: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winsuggestion: Exercise the oversized-bin execution path.
Both calls pass
nullptr, so the dispatch performs no work. Allocatenum_bins_overflowcounters and runHistogramEvenwith non-null temporary storage, especially forint16_t.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: fecbe981-f97a-430c-9d46-c2700be6813b
📒 Files selected for processing (11)
cub/benchmarks/bench/histogram/even.cucub/benchmarks/bench/histogram/histogram_common.cuhcub/benchmarks/bench/histogram/histogram_inputs.cuhcub/benchmarks/bench/histogram/multi/even.cucub/benchmarks/bench/histogram/multi/range.cucub/benchmarks/bench/histogram/range.cucub/cub/device/dispatch/dispatch_histogram.cuhcub/cub/device/dispatch/kernels/kernel_histogram.cuhcub/test/catch2_test_device_histogram.cucub/test/catch2_test_device_histogram_input_shapes.cucub/test/catch2_test_device_histogram_thread_local_cache.cu
🚧 Files skipped from review as they are similar to previous changes (7)
- cub/test/catch2_test_device_histogram_thread_local_cache.cu
- cub/cub/device/dispatch/kernels/kernel_histogram.cuh
- cub/benchmarks/bench/histogram/multi/range.cu
- cub/benchmarks/bench/histogram/multi/even.cu
- cub/benchmarks/bench/histogram/histogram_inputs.cuh
- cub/cub/device/dispatch/dispatch_histogram.cuh
- cub/benchmarks/bench/histogram/histogram_common.cuh
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) <noreply@anthropic.com>
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) <noreply@anthropic.com>
`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) <noreply@anthropic.com>
(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) <noreply@anthropic.com> (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) <noreply@anthropic.com>
(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) <noreply@anthropic.com>
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) <noreply@anthropic.com>
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) <noreply@anthropic.com>
… 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) <noreply@anthropic.com>
robobryce
force-pushed
the
pr/histocache/benchmark-testing
branch
from
August 30, 2026 20:06
8cf15bc to
c974df6
Compare
Contributor
|
Note GitHub couldn't provide a complete incremental comparison for this pull request, so CodeRabbit is performing a full review instead. This review may take a little longer. |
Contributor
There was a problem hiding this comment.
Actionable comments posted: 2
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: bcdcca4f-71fd-4562-a260-e9db68f2e0a0
📒 Files selected for processing (11)
cub/benchmarks/bench/histogram/even.cucub/benchmarks/bench/histogram/histogram_common.cuhcub/benchmarks/bench/histogram/histogram_inputs.cuhcub/benchmarks/bench/histogram/multi/even.cucub/benchmarks/bench/histogram/multi/range.cucub/benchmarks/bench/histogram/range.cucub/cub/device/dispatch/dispatch_histogram.cuhcub/cub/device/dispatch/kernels/kernel_histogram.cuhcub/test/catch2_test_device_histogram.cucub/test/catch2_test_device_histogram_input_shapes.cucub/test/catch2_test_device_histogram_thread_local_cache.cu
🚧 Files skipped from review as they are similar to previous changes (10)
- cub/cub/device/dispatch/dispatch_histogram.cuh
- cub/test/catch2_test_device_histogram_thread_local_cache.cu
- cub/cub/device/dispatch/kernels/kernel_histogram.cuh
- cub/benchmarks/bench/histogram/multi/range.cu
- cub/benchmarks/bench/histogram/multi/even.cu
- cub/benchmarks/bench/histogram/even.cu
- cub/test/catch2_test_device_histogram.cu
- cub/benchmarks/bench/histogram/histogram_inputs.cuh
- cub/benchmarks/bench/histogram/range.cu
- cub/benchmarks/bench/histogram/histogram_common.cuh
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
Contributor
|
/ok to test f80e851 |
Contributor
🔬 CUB benchmark SASS comparisonHow to request a benchmark run
Targets with a SASS change
|
Contributor
🥳 CI Workflow Results🟩 Finished in 1h 29m: Pass: 100%/284 | Total: 2d 22h | Max: 56m 15s | Hits: 85%/251795See results here. AI failure analysis1. WholeGraph GoogleTest fetch returned HTTP 504 · 1 jobExplanation: WholeGraph configuration failed before compilation because CMake FetchContent received an HTTP 504 while downloading GoogleTest from GitHub. The PR changes only CUB histogram code, and the earlier tests-disabled WholeGraph build succeeded, so this is an external dependency-fetch failure rather than evidence of a product-code regression. Evidence: Copy this prompt into a coding agentJobs: |
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
2 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.
Why
DeviceHistogram tuning needs benchmark inputs that exercise different distributions and ordering patterns, plus an independent way to reject incorrect results. The previous benchmark exposed only a narrow entropy-style input family and did not provide an opt-in bin-by-bin reference check.
While adding that coverage, the new cases exposed integer-overflow bugs in evenly spaced bin computation and invalid benchmark cells whose sample count or multi-channel row stride did not fit the selected offset type.
What changed
Deterministic input shapes
The single- and multi-channel EVEN and RANGE benchmarks now share deterministic input generators for:
The
hash_synonymandstale_residentnames remain reserved as explicit placeholders. Their generators intentionally report that they are unavailable until the shared-memory cache policy lands, avoiding duplicated cache-size constants in this benchmark-only PR.The default matrix remains close to the previous runtime: 216 configurations per benchmark instead of 192. It uses element counts of 65,536, 4,000,000, and 67,000,000; bin counts of 33, 2,048, 16,384, and 2,000,003; and three representative shapes covering uniform, concentrated, and ordered traffic. The non-power-of-two values reduce tuning bias, while the lower top element count keeps the four-channel, 64-bit sample case to roughly 2 GiB of input.
RANGE benchmarks generate reproducible non-uniform integer levels so they exercise the arbitrary-level search path rather than collapsing back to uniform spacing. Level construction now detects exhausted integral ranges before applying a monotonicity repair.
Opt-in correctness verification
Set
CUB_BENCH_HISTOGRAM_VERIFY=1(also acceptstrue,yes, oron) to run one untimed invocation and compare every output bin with an independent device reference. Correctness checks remain off by default, matching the existing benchmark behavior. A mismatch terminates the benchmark process instead of being reported as a skipped measurement.The benchmark counter and offset types can be overridden for dedicated 64-bit builds. The reference path supports signed 64-bit counters, and benchmark cells are skipped when their sample count or multi-channel row stride cannot be represented by
OffsetT.DeviceHistogram fixes and tests
The integer EVEN transform stores scaling operands in the widened arithmetic type and computes signed level differences through their same-width unsigned representation. This avoids undefined signed subtraction and preserves level ranges wider than
LevelTwithout changing the signed public API or its range comparisons. Boolean samples retain their previous pass-through behavior without instantiatingmake_unsigned_t<bool>.Tests cover the input-shape contract, negative minimum levels, bin counts wider than
LevelT, non-uniform RANGE fixtures, and the existing thread-local RANGE-detection cache across streams, host threads, and device changes. The cache tests explicitly preserve CUDA device affinity and restore the caller's device.Scope and relationship to other work
This PR contains benchmark infrastructure, focused DeviceHistogram arithmetic fixes, and tests. It does not contain the shared-memory privatization algorithms in #10556 or the cooperative high-bin cache work in #10568. Those PRs are independently based on
main; this PR is not a stack base for either one.The exploratory Python sweep and plotting scripts were removed from CCCL and moved to
robobryce/histocache-benchmarking-scripts, as requested in review. The implementation originated from the raw research snapshots in #10547 and #10554, but scratch artifacts and experimental histogram algorithms are not included here.The repeated benchmark-local persisting-L2 reset was also removed. NVIDIA/nvbench#459 implements the corresponding per-benchmark NVBench option, but it is still a draft whose current commit is hosted only in a contributor fork. This PR therefore does not pin all of CCCL to that fork; the four histogram registrations can enable the NVBench option after #459 merges and CCCL advances its NVBench pin.
Validation
Formatting and repository checks:
Fresh CUDA 13.3.33 / GCC 13.3 / SM100 test build after the August 30 rebase:
Built successfully:
cub.test.device.histogramcub.test.device.histogram_input_shapescub.test.device.histogram_thread_local_cachePassed on an NVIDIA B200 after the final rebase onto
main:cub.test.device.histogram.lid_0cub.test.device.histogram.lid_1cub.test.device.histogram.lid_2cub.test.device.histogram_input_shapes.lid_0cub.test.device.histogram_thread_local_cache.lid_0Fresh benchmark build with CMake 4.3.2, CUDA 13.3.33, GCC 13.3, and SM100 after the August 30 rebase:
Built successfully:
cub.bench.histogram.even.basecub.bench.histogram.range.basecub.bench.histogram.multi.even.basecub.bench.histogram.multi.range.baseAfter the final rebase, each benchmark binary passed a correctness-enabled smoke run with
SampleT=I32,CounterT=I32,OffsetT=I32, 65,536 elements, 2,048 bins, and thestrided_sweepinput shape.