[CUB] Fix HistogramEven bin arithmetic for level ranges wider than the common type - #10996
[CUB] Fix HistogramEven bin arithmetic for level ranges wider than the common type#10996VaggelisGian wants to merge 2 commits into
Conversation
…e common type The even-bin computation subtracted levels in the narrow common type of LevelT and SampleT before widening to the arithmetic type, so a range like full int16 wrapped to a negative span (silently binning every sample into bin 0), wide int32 ranges overflowed and were spuriously rejected with cudaErrorInvalidValue despite being far below the documented uint64_t product bound, and the num_bins representability check from NVIDIA#6908 only fired for some counts because the count was narrowed through CommonT first. Store the integral scale fraction in the widened arithmetic type, compute both differences after widening their operands, reject bin counts above the common type maximum explicitly instead of relying on narrowing wrap, and pass the raw bin count from all dispatch sites. The custom-type and __int128 path keeps its previous semantics by casting the fraction fields back to the common type. Test Plan: NVIDIA#10975 reproducers, CUDA 12.8, sm_90: int16 full-range 100 bins: bins {0,50,99} counted (was all in bin 0) int32 [-1.5e9,1.5e9) 100 bins: cudaSuccess (was cudaErrorInvalidValue) ci/util/build_and_test_targets.sh --preset cub-cpp20 --build-targets cub.test.device.histogram --ctest-targets '^cub[.]test[.]device[.]histogram$' -> Passed (3m24s), including the existing rejection-contract tests and new regression tests for both repro configurations
|
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 (1)
🚧 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; 11 remain after this review. 📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesThe integral HistogramEven integral arithmetic
Assessment against linked issues
Suggested reviewers: Merge Risk: ⚪ Minimal · up to This change corrects histogram bin arithmetic for wider integral level ranges and adds regression coverage; no actionable merge-blocking risk remains beyond normal checks and review. Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
🧹 Nitpick comments (1)
cub/test/catch2_test_device_histogram.cu (1)
841-851: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Declare immutable test fixtures
const, and declareh_samplesconstexprin both tests. Keepd_histogrammutable becausehistogram_evenwrites the output buffer. As per coding guidelines, “All variables that are not modified must be declaredconst” and “All variables that can be evaluated at compile time must be declaredconstexpr.”Also applies to: 863-873
Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: e0feea89-d5ed-454c-90e0-d4d359cedc5c
📒 Files selected for processing (3)
cub/cub/device/dispatch/dispatch_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; 10 remain after this review.
Per coding guidelines: variables evaluable at compile time are declared constexpr.
|
/ok to test 2f65cb1 |
1 participant
Description
closes #10975
cub::DeviceHistogram::HistogramEvenmishandled integral level ranges wider than the common type ofLevelTandSampleT. Three symptoms, all from the same root cause: the bin arithmetic subtracted levels in the narrow common type before widening.int16_tlevels over [-32768, 32767) put every sample in bin 0 (max_level - min_levelwrapped to -1 inint16_t), returningcudaSuccess.int32_trange with 100 bins was rejected withcudaErrorInvalidValueeven though(upper - lower) * (num_levels - 1)= 3e11 is far below the documenteduint64_tbound; the subtraction overflowedint32_tfirst.CommonTslipped past the Fix overflow issue in histogram even benchmark when the number of bins exceeds whatLevelTcan represent #6908 rejection for some values because the count was narrowed throughstatic_cast<CommonT>before the check and wrapped into a huge unsigned number there.Changes
ScaleTransformstores the integral scale fraction inIntArithmeticT(the widened arithmetic type) instead ofCommonT, computesrangeas a widened subtraction, and computessample - min_levelafter widening both operands. Modular unsigned subtraction of two cast operands is exact here because samples belowmin_levelare rejected before bin computation.__int128bin path casts the fraction fields back toCommonT, keeping its semantics unchanged.MayOverflownow rejects bin counts aboveCommonT's maximum explicitly (same observable contract as Fix overflow issue in histogram even benchmark when the number of bins exceeds whatLevelTcan represent #6908, without relying on narrowing wrap), then checks the widened range againstIntArithmeticTmax divided by the bin count. Both host-init call sites pass the raw bin count instead of pre-narrowing it; the device-init paths get the same rejection they previously lacked.Documented contract preserved: rejection iff
(upper_level[i] - lower_level[i]) * (num_levels[i] - 1)exceeds the widened arithmetic type (128-bit when the common type is 128 bits wide). One deliberate gap remains: small sample/level types use auint32_tfast arithmetic path, so products between 2^32 and 2^64 are still rejected where the docs promise acceptance. Fixing that would move those configs to 64-bit division in the per-sample hot loop; I left it alone pending maintainer input on whether the docs or the current policy should change.Verification
Reproducers from the issue, CUDA 12.8, sm_90: full-range int16 now reports bins {0, 50, 99} for samples {-32768, 0, 32766} instead of all-in-bin-0; the int32 case returns success instead of
cudaErrorInvalidValue. New regression tests cover both configurations. The existing "num_bins exceeds LevelT range" and "bin computation does not overflow" tests pin the rejection contract and still pass.Checklist