[CUB] GPU-to-GPU DeviceReduce with device resident problem size#9740
[CUB] GPU-to-GPU DeviceReduce with device resident problem size#9740NaderAlAwar wants to merge 29 commits into
DeviceReduce with device resident problem size#9740Conversation
…ferred-gpu-to-gpu-num-items
…ferred-gpu-to-gpu-num-items
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
|
/ok to test 52aa029 |
This comment has been minimized.
This comment has been minimized.
…ferred-gpu-to-gpu-num-items
|
/ok to test 4808d52 |
|
/ok to test 40ee0fc |
This comment has been minimized.
This comment has been minimized.
|
@coderabbitai full review |
✅ Action performedFull review finished. |
|
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 (2)
📝 WalkthroughSummary by CodeRabbit
WalkthroughAdds device-resident ChangesDeferred deterministic reduce support
Assessment against linked issues
Possibly related PRs
Suggested reviewers: 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: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 6bb7f349-d5b1-402e-af08-6c5790279f17
📒 Files selected for processing (7)
cub/benchmarks/bench/reduce/deferred_deterministic.cucub/benchmarks/bench/reduce/deterministic.cucub/cub/device/device_reduce.cuhcub/cub/device/dispatch/dispatch_reduce_deterministic.cuhcub/cub/device/dispatch/kernels/kernel_reduce_deterministic.cuhcub/test/catch2_test_device_reduce_deferred.cucub/test/catch2_test_device_reduce_env.cu
This comment has been minimized.
This comment has been minimized.
fbusato
left a comment
fbusato
left a comment
There was a problem hiding this comment.
reviewed only the cub/device code for now
| { | ||
| using element_t = it_value_t<KernelNumItemsT>; | ||
| using type = ::cuda::std::conditional_t< | ||
| ::cuda::std::is_signed_v<element_t> && sizeof(element_t) == 4, |
There was a problem hiding this comment.
question. could we use is_same_v<int, element_t> instead?
There was a problem hiding this comment.
Likely not because on windows we have long which is also 32 bits, so it should go through fast path
| typename TransformOpT = ::cuda::std::identity> | ||
| _CCCL_KERNEL_ATTRIBUTES __launch_bounds__( | ||
| int(current_policy<PolicySelector>().single_tile.threads_per_block), | ||
| 1) void DeterministicDeviceReduceDeferredSingleTileKernel(_CCCL_GRID_CONSTANT const InputIteratorT d_in, |
There was a problem hiding this comment.
considering the recent regressions, I would refrain from using _CCCL_GRID_CONSTANT
There was a problem hiding this comment.
Since all of our kernels have it, I would prefer to keep it for consistency. This also helps if it turns out to be a compiler bug that gets fixed later
🥳 CI Workflow Results🟩 Finished in 2h 39m: Pass: 100%/287 | Total: 12d 17h | Max: 2h 38m | Hits: 20%/975031See results here. |
2 participants
Description
closes #9730
merge first #9722
This extends device resident problem size to GPU-to-GPU determinism.
No SASS changes to any existing benchmarks.
We expect the device-resident problem size to be slightly slower: the host cannot read the problem size, so it must assume it is large, launching a worst-case grid and always taking the two-pass path (there is no single-tile shortcut for small problems). The kernels read the problem size on device and trim the worst-case grid to the size the host would have computed: surplus blocks exit early, and the second pass derives the same block count to consume exactly the partials that were produced. Because the RFA accumulator is partition-independent, results are bitwise identical to the host-provided num_items path (this is verified by the tests, including problem sizes above INT32_MAX, where the two paths chunk the input differently). Compared to the host-provided num_items, the results show an overhead of at most ~7 us, which is most noticeable at the smallest problem sizes, and low single-digit percentages elsewhere.
When the problem-size element is anything other than a signed 32-bit integer we observe larger overheads at mid-to-large sizes (an unsigned 32-bit element cannot guarantee the problem fits a single chunk, so it is widened and takes the same path as 64-bit elements). Unlike the host-provided path, which chunks the input on the host so the kernel only ever uses 32-bit index arithmetic, the deferred kernel performs the equivalent chunking on device. The hot loop remains 32-bit, but carrying the wide problem size across the loop widens the kernel's register footprint relative to the signed-32-bit instantiation (F32: 40 vs. 32 registers, F64: 44 vs. 40 on SM90), which reduces occupancy. Profiling with Nsight Compute confirms this is the cause: achieved occupancy drops from 76% to 56% at F32/2^24, where the kernel is latency-bound, and the effect fades as bandwidth saturates at larger sizes. Users whose problem sizes fit in 31 bits should prefer a signed 32-bit element for the deferred source; a future improvement could use cuda::args::bounds to statically select the fast kernel from a bounded wider source.
GPU-to-GPU reduce
Checklist