[CUB] Fix silent sample loss in DeviceHistogram for signed bytes and wide histograms - #10993
[CUB] Fix silent sample loss in DeviceHistogram for signed bytes and wide histograms#10993VaggelisGian wants to merge 2 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
|
/ok to test |
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (2)
🚧 Files skipped from review as they are similar to previous changes (1)
Included review availability: Your plan provides up to 12 included reviews per hour; 9 remain after this review. 📝 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: 🔵 Low · up to The PR fixes histogram sample loss and count corruption, with the remaining risk limited to a localized naming-convention issue in the implementation; it is mergeable with owner awareness or a follow-up cleanup. Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
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.
|
/ok to test 7cffd55 |
1 participant
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