[CUB] Add fast constant division utility by robobryce · Pull Request #11304 · NVIDIA/cccl

robobryce · 2026-09-09T16:34:27Z

Why

The cooperative histogram work in #10568 introduced a small precomputed divider so an invariant unsigned divisor does not require a hardware integer division for every sample. Review requested that this generally useful facility live in CUB's internal utilities and be reviewed independently of both #10568 and the benchmark work in #10555.

Keeping the divider separate also gives its arithmetic edge cases direct coverage. In particular, divisors above half of the integer range cannot use the normal magic-number setup because constructing its numerator would require shifting by the full width of the widened type.

What changed

Added cub::detail::fast_divide_by_constant to the existing cub/detail/fast_modulo_division.cuh internal utility header.
Preserved the original safe identity behavior for divisors zero and one.
Use a right shift for powers of two.
Precompute a multiplier and shift for other representable 32-bit and 64-bit unsigned divisors, then use multiply-high when dividing.
Fall back to ordinary division for full-width non-power-of-two divisors, avoiding an undefined full-width shift during precomputation.
Use cuda::std::countl_zero and fixed-width cuda::std integer types. The 64-bit host path also retains a portable limb-based multiply-high and precomputation fallback for toolchains without 128-bit integers.

For example, repeatedly dividing values by 3, 10, or 2^63 - 1 uses precomputed state, while a divisor such as 2^63 + 1 safely selects the hardware fallback. Divisors 0 and 1 return the numerator unchanged, matching the extracted facility's disabled/identity state.

Tests

Added focused host and device tests for uint32_t and uint64_t. They compare every result against the corresponding exact quotient across:

zero and one;
small and top-bit powers of two;
representative non-powers of two;
values immediately below and above the half-range boundary;
maximum-width divisors, including max - 1 and max;
numerator boundaries around each divisor and its largest representable multiple;
4,096 deterministic full-width numerator values per type.

Validation was run after rebasing onto current upstream/main (4b8e07fac3f8acc78712ee14da73eed3c99338c4):

CUB C++17, CUDA 13.3.33, GCC 13.3, sm100-real,
CCCL_ENABLE_WERROR=ON, CCCL_ENABLE_PRAGMA_SYSTEM_HEADER=OFF

cub.test.fast_divide_by_constant: passed
cub.test.fast_div_mod: passed

CUB C++20 with CCCL_DISABLE_INT128_SUPPORT, CUDA 13.3.33, GCC 13.3,
sm100-real, CCCL_ENABLE_WERROR=ON, CCCL_ENABLE_PRAGMA_SYSTEM_HEADER=OFF

cub.test.fast_divide_by_constant: passed
cub.test.fast_div_mod: passed

pre-commit run --files cub/cub/detail/fast_modulo_division.cuh cub/test/internal/catch2_test_fast_divide_by_constant.cu also passed.

Scope

This PR intentionally does not modify DeviceHistogram implementation, dispatch, tuning, benchmarks, or tests. PR #10568 is removing its local copy of this optimization and does not depend on this PR. Any future algorithm adoption of this utility will be reviewed separately; this branch is not stacked on either #10568 or #10555.

copy-pr-bot · 2026-09-09T16:34:30Z

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.

fbusato

We already have it cuda::fast_mod_div https://nvidia.github.io/cccl/unstable/libcudacxx/extended_api/math/fast_mod_div.html#libcudacxx-extended-api-math-fast-mod-div

coderabbitai · 2026-09-09T16:42:13Z

📝 Summary

Summary by CodeRabbit

Performance Improvements
- Improved unsigned integer division for 32-bit and 64-bit values.
- Added optimized handling for common constant divisors, including zero, one, and powers of two.
- Supports consistent execution across host and device environments, helping arithmetic-heavy workloads run more efficiently.
Reliability
- Expanded validation across boundary values, large operands, and randomized inputs to improve confidence in division results.

Walkthrough

Adds fast_divide_by_constant for unsigned 32-bit and 64-bit host/device division. The implementation selects optimized division strategies and includes 64-bit multiply-high support. New tests validate host and CUDA results across representative divisors and numerators.

Changes

Constant Division

Layer / File(s)	Summary
Fast divider implementation `cub/cub/detail/fast_modulo_division.cuh`	Adds `fast_divide_by_constant` with identity, power-of-two, hardware, and multiply-shift division paths. Supports host/device initialization and division for unsigned 32-bit and 64-bit values.
Host and device validation `cub/test/internal/catch2_test_fast_divide_by_constant.cu`	Tests zero-safe reference results, boundary cases, representative numerators, host division, CUDA kernel execution, synchronization, and errors.

Suggested reviewers: bernhardmgruber, davebayer

Priority: ⬇️ Low

Merge Risk: 🔵 Low · up to 8a37b

The divider is mergeable, but large divisors may unnecessarily reduce CUDA performance by using full-width division where a comparison would suffice. The duplication concern is maintainability follow-up rather than a current runtime blocker.

_{Comment @coderabbitai help to get the list of available commands.}

coderabbitai

🧹 Nitpick comments (2)

cub/cub/detail/fast_modulo_division.cuh (2)
206-212: 🚀 Performance & Scalability | 🔵 Trivial | ⚡ Quick win

suggestion: replace mode::hardware with a comparison.

log2_divisor == bits implies divisor > 2^(bits-1), so the quotient is only 0 or 1. Line 259 then issues a full-width integer division, which is expensive on device for 64-bit operands. fast_div_mod already handles this case with a comparison (lines 343-347).
     identity,
     shift,
     multiply_shift,
-    hardware
+    compare
   };
     const int log2_divisor = ceil_log2(divisor);
     if (log2_divisor == bits)
     {
       magic_ = divisor;
       shift_ = 0;
-      mode_  = mode::hardware;
+      mode_  = mode::compare;
       return;
     }
-    if (mode_ == mode::hardware)
+    if (mode_ == mode::compare)
     {
-      return numerator / magic_;
+      return numerator >= magic_ ? UInt{1} : UInt{0};
     }
Source: Path instructions

122-126: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚖️ Poor tradeoff

suggestion: the header now contains two independent magic-multiplier dividers.

fast_divide_by_constant and fast_div_mod duplicate divisor classification, magic computation, and multiply-high logic. Consider building fast_div_mod on top of the new class, or document why the two must stay separate.

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 35981c9c-0bd1-4d07-9c54-c4f78cb360a6

📥 Commits

Reviewing files that changed from the base of the PR and between 4b8e07f and 8a37b3c.

📒 Files selected for processing (2)

cub/cub/detail/fast_modulo_division.cuh
cub/test/internal/catch2_test_fast_divide_by_constant.cu

Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.

davebayer · 2026-09-10T06:10:59Z

+  [[nodiscard]] _CCCL_HOST_DEVICE _CCCL_FORCEINLINE static int ceil_log2(UInt divisor) noexcept
+  {
+    return divisor <= UInt{1} ? 0 : bits - ::cuda::std::countl_zero(divisor - UInt{1});
+  }


We have cuda::ceil_ilog2

davebayer · 2026-09-10T06:11:09Z

+    return divisor <= UInt{1} ? 0 : bits - ::cuda::std::countl_zero(divisor - UInt{1});
+  }
+
+  [[nodiscard]] _CCCL_HOST_DEVICE _CCCL_FORCEINLINE static UInt multiply_high(UInt lhs, UInt rhs) noexcept


We have cuda::mul_hi

davebayer · 2026-09-10T06:12:25Z

+      mode_  = mode::identity;
+      return;
+    }
+    if ((divisor & (divisor - UInt{1})) == UInt{0})


Suggested change

if ((divisor & (divisor - UInt{1})) == UInt{0})

if (cuda::is_power_of_two(divisor))

[cub] Add fast constant division utility

8a37b3c

robobryce requested a review from a team as a code owner September 9, 2026 16:34

robobryce requested a review from NaderAlAwar September 9, 2026 16:34

github-project-automation Bot added this to CCCL Sep 9, 2026

github-project-automation Bot moved this to Todo in CCCL Sep 9, 2026

cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Sep 9, 2026

fbusato requested changes Sep 9, 2026

View reviewed changes

github-project-automation Bot moved this from In Review to In Progress in CCCL Sep 9, 2026

coderabbitai Bot reviewed Sep 9, 2026

View reviewed changes

davebayer reviewed Sep 10, 2026

View reviewed changes

Navigation Menu

Uh oh!

[CUB] Add fast constant division utility - #11304

[CUB] Add fast constant division utility#11304
robobryce wants to merge 1 commit into
NVIDIA:mainfrom
robobryce:pr/histocache/fast-divide

robobryce commented Sep 9, 2026 •

edited

Loading

Uh oh!

copy-pr-bot Bot commented Sep 9, 2026

Uh oh!

fbusato left a comment

Uh oh!

coderabbitai Bot commented Sep 9, 2026 •

edited

Loading

Summary by CodeRabbit

Walkthrough

Changes

Uh oh!

coderabbitai Bot left a comment

Uh oh!

davebayer Sep 10, 2026

Uh oh!

davebayer Sep 10, 2026

Uh oh!

davebayer Sep 10, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

	if ((divisor & (divisor - UInt{1})) == UInt{0})
	if (cuda::is_power_of_two(divisor))

Uh oh!

Conversation

robobryce commented Sep 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Why

What changed

Tests

Scope

Uh oh!

copy-pr-bot Bot commented Sep 9, 2026

Uh oh!

fbusato left a comment

Choose a reason for hiding this comment

Uh oh!

coderabbitai Bot commented Sep 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary by CodeRabbit

Walkthrough

Changes

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

davebayer Sep 10, 2026

Choose a reason for hiding this comment

Uh oh!

davebayer Sep 10, 2026

Choose a reason for hiding this comment

Uh oh!

davebayer Sep 10, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

robobryce commented Sep 9, 2026 •

edited

Loading

coderabbitai Bot commented Sep 9, 2026 •

edited

Loading