[CUB] Fix silent sample loss in DeviceHistogram for signed bytes and wide histograms - #10993
Open
VaggelisGian wants to merge 3 commits into
Open
[CUB] Fix silent sample loss in DeviceHistogram for signed bytes and wide histograms#10993VaggelisGian wants to merge 3 commits into
VaggelisGian wants to merge 3 commits into
Conversation
Byte-sized samples took a fast path that indexes a 256-bin privatized histogram with the raw sample value, so negative int8_t samples produced negative bin indices and were never counted (NVIDIA#10977). Separately, when folding privatized counts into the output histogram, each bin index was cast through the sample type before the output decode op, so indices that are not exactly representable in the sample type were altered on the round trip: __nv_bfloat16 merged bins above 256 and could write one element past the output histogram, __half broke above 2048 bins, and int16_t dropped every bin >= 32768 (NVIDIA#10976). Restrict the byte-sample fast path to unsigned 8-bit sample types so signed bytes use the generic ScaleTransform/SearchTransform path, pass privatized bin indices to the output decode op as int instead of round-tripping them through the sample type, and apply the same signedness gate in c/parallel's byte-sample selection. Test Plan: Reproducers from both issues, CUDA 12.8, sm_90: full-range int8_t 256-bin histogram: 256 of 256 counted (was 128) bf16 HistogramRange 300 consecutive levels: 0 wrong bins (was 33 + OOB atomicAdd) ci/util/build_and_test_targets.sh --preset cub-cpp20 --build-targets cub.test.device.histogram --ctest-targets '^cub[.]test[.]device[.]histogram$' -> Passed (5m28s), includes new regression tests for negative int8 samples, the full 8-bit domain, 65536-bin int16, and half > 2048 bins neighbor targets cub.test.device.histogram_api / _env_api / _custom_policy_hub / cub.test.block.histogram compile clean
Contributor
Contributor
Author
|
/ok to test |
Contributor
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesSigned 8-bit samples now use the non-byte histogram path. Privatized bin indices no longer convert through the sample type. Regression tests cover negative ChangesHistogram correctness
Assessment against linked issues
Suggested reviewers: Merge Risk: ⚪ Minimal · up to This PR fixes silent histogram sample loss and count corruption for signed bytes and wide histograms. No actionable merge-blocking risk remains; a localized naming-style follow-up does not affect runtime correctness. Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Actionable comments posted: 2
🧹 Nitpick comments (1)
cub/test/catch2_test_device_histogram.cu (1)
843-844: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: Declare immutable test-state variables
const.const auto*does not maked_sample_ptrimmutable. Useconst auto* const d_sample_ptrand declare theh_expectedvectorsconst. As per coding guidelines, “All variables that are not modified must be declaredconst.”Also applies to: 884-885, 933-935, 1001-1001
Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 52a18ca5-b8b6-40b6-9681-e2d5bdb72c2a
📒 Files selected for processing (5)
c/parallel/src/histogram.cucub/cub/agent/agent_histogram.cuhcub/cub/device/device_histogram.cuhcub/cub/device/dispatch/kernels/kernel_histogram.cuhcub/test/catch2_test_device_histogram.cu
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
Make the dispatch-time bool_constant objects constexpr and the test sample pointers fully const.
Contributor
Author
|
/ok to test 7cffd55 |
griwes
approved these changes
Aug 28, 2026
griwes
left a comment
griwes
left a comment
Contributor
There was a problem hiding this comment.
Reviewed throughout. This looks good.
Contributor
|
/ok to test 7cffd55 |
Contributor
|
@VaggelisGian copy-pr-bot only listens to maintainers :P |
Contributor
🔬 CUB benchmark SASS comparisonHow to request a benchmark run
Targets with a SASS change
|
Contributor
😬 CI Workflow Results🟥 Finished in 1h 30m: Pass: 96%/276 | Total: 2d 22h | Max: 1h 10m | Hits: 84%/222822See results here. AI failure analysis1. CUB DeviceHistogram byte-sample tag fails GCC 7 with unused-but-set-variable · 8 jobsExplanation: The change introduced initialized constexpr tag objects in four DeviceHistogram overloads. NVCC's GCC 7 host compilation diagnoses these locals as set but unused, and -Werror aborts all eight CTK and launch-mode builds. Evidence: Copy this prompt into a coding agentJobs:
2. CUB histogram half regression test fails clang-tidy narrowing conversion · 1 jobExplanation: The newly added `exact_half` lambda implicitly converts the integer `k` to float during division. Clang-tidy treats that conversion as a narrowing warning, which is promoted to an error. Evidence: Copy this prompt into a coding agentJobs: |
Replace the constexpr bool_constant locals in the four DeviceHistogram overloads with a type alias and pass a temporary to dispatch, avoiding -Wunused-but-set-variable on the NVCC GCC7 host pass. Explicitly convert the integer level index to float in the half regression test to satisfy clang-tidy narrowing checks.
Contributor
Author
|
Thanks for running the CI, understood on the vetting. Run 33213561004 surfaced two compile diagnostics from this PR, both fixed in a7f211c:
On the SASS diff for the 4 histogram benchmark targets: that is expected, the bin arithmetic fix changes the generated code. Happy to request a benchmark run if you want the perf numbers. |
Contributor
There was a problem hiding this comment.
🧹 Nitpick comments (1)
cub/test/catch2_test_device_histogram.cu (1)
977-977: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Use uniform initialization for this
half_tconstruction:half_t{static_cast<float>(k) / 1024.0f}.Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 964737f0-733e-4348-b3e8-567b31f7cae9
📒 Files selected for processing (2)
cub/cub/device/device_histogram.cuhcub/test/catch2_test_device_histogram.cu
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
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.
Description
closes #10977
closes #10976
cub::DeviceHistogramsilently lost samples and corrupted counts in two related ways, both returningcudaSuccess:Negative signed byte samples vanished ([BUG]: cub::DeviceHistogram silently drops all negative int8_t samples (Even and Range) #10977). Byte-sized sample types take a fast path that uses the raw sample value as a privatized bin index into a 256-bin shared-memory histogram (
PassThruTransform). A negativeint8_tsample produced a negative bin index, which the agent treats as "outside the histogram", so every negative sample was dropped.Bin indices were rounded or wrapped after binning ([BUG]: [BUG]: cub::DeviceHistogram mis-counts bins and writes out of bounds when the number of bins exceeds what the sample type can represent exactly #10976). When folding per-block privatized histograms into the output,
agent_histogram.cuhcast each privatized bin index back through the sample type before handing it to the output decode op. Any bin index not exactly representable in the sample type got altered on that round trip:__nv_bfloat16above 256 bins merged neighboring bins and could write one element past the end of the output histogram,__halfbroke past 2048 bins,int16_tdropped every bin >= 32768, andfloatbreaks past 2^24 bins. This affectedHistogramEven,HistogramRange, and theMultivariants.The two bugs compound: fixing only #10977 by routing signed bytes to the generic path still drops bins 128..255 for full-range
int8_thistograms, because those bin indices wrap negative in the same round trip.Changes
sizeof(SampleT) == 1 && !is_signed_v<SampleT>at all four public entry points. Signed byte types take the genericScaleTransform/SearchTransformpath, which handles negative samples and levels correctly.ScaleTransform::BinSelect's sample parameter and pass the privatized bin index fromStoreOutputasint. The decode ops are exact for any representable bin index; unsigned byte paths are unaffected because their indices are already <= 255.c/parallel/src/histogram.cu, which picked its byte-sample specialization purely fromvalue_type.size == 1.Unsigned byte samples keep the fast path with unchanged behavior. Plain
charis signed on theplatforms we test (and on MSVC by default), so existing
char-sampled histograms now take the genericpath: correctness improves wherever levels go negative, and the trade-off is a performance-only change
for those configs.
Verification
Reproducers from the two issues, CUDA 12.8, sm_90:
int8_thistogram (256 bins over [-128, 128)): 128 of 256 counted before, 256 of 256 after.HistogramRangewith 300 consecutive bf16 levels, one sample per bin: 33 wrong bins before (plus an out-of-bounds atomicAdd), 0 wrong bins after.New regression tests in
catch2_test_device_histogram.cu: negative int8 samples over [-60, 64), the full 8-bit domain with 256 bins (Even and Range), 65536-bin int16 histograms (Even and Range), and a guarded__halfcase with 3000 exactly representable bins. The bf16 configuration is covered by the reproducer above; adding it to the catch2 test needs bf16 support plumbing this test file does not have yet.Note for reviewers: if the performance of signed-byte histograms matters, a follow-up could teach the pass-thru path to shift signed samples into [0, 256) together with a matching offset in the output decode op; I kept the fix minimal instead.
Checklist