NVIDIA / NVIDIA/cccl

[BUG]: cub::DeviceReduce gpu_to_gpu determinism writes out of bounds when num_items is an exact multiple of INT32_MAX

Open
#9,771 0 comments 1 reaction 1 assignee Claimed by @NaderAlAwar View on GitHub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.