[BUG]: cub::DeviceReduce gpu_to_gpu determinism writes out of bounds when num_items is an exact multiple of INT32_MAX
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
### Is this a duplicate?
- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Type of Bug
Runtime Error
### Component
CUB
### Describe the bug
The deterministic (RFA / `gpu_to_gpu`) reduce splits large problems into host-side chunks of `INT32_MAX` items (`detail::rfa::invoke_passes` in `cub/cub/device/dispatch/dispatch_reduce_deterministic.cuh`). The temporary-storage sizing mishandles problem sizes that are an exact multiple of the chunk size: the final (full) chunk contributes zero blocks to the partials count, so the temporary storage is under-allocated while the launch loop still runs the final chunk at full grid size, writing its per-block partials past the end of the allocation. The second pass then also consumes the wrong number of partials.
### How to Reproduce
```cpp
const auto in = cuda::constant_iterator{1.0};
const auto env = cuda::execution::require(cuda::execution::determinism::gpu_to_gpu);
// num_items = INT32_MAX (and 2 * INT32_MAX as int64) -> wrong result; OOB visible under compute-sanitizer
cub::DeviceReduce::Reduce(in, out.begin(), std::int64_t{INT32_MAX}, cuda::std::plus<>{}, 0.0, env);
```
Expected `INT32_MAX` (exact in double); observed `init` (for k == 1) or a corrupted/incorrect sum
(k >= 2). `compute-sanitizer --tool memcheck` flags the out-of-bounds stores.
### Expected behavior
Size the last chunk from the actual final-chunk item count instead of the remainder:
```cpp
const int last_chunk_num_items = has_partial_chunk ? partial_chunk_size : num_items_per_chunk;
const int last_chunk_tile_grid_size = ::cuda::ceil_div(last_chunk_num_items, reduce_config.tile_size);
```
### Reproduction link
_No response_
### Operating System
_No response_
### nvidia-smi output
_No response_
### NVCC version
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.