[EPIC]: Thrust/CUB shouldn't invoke user-defined operators on out-of-bounds data
- 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 request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Area
General CCCL
### Is your feature request related to a problem? Please describe.
As a user of Thrust/CUB, I can provide a custom operator to the many algorithms. My operator may only be defined for the values that are present in the input. As a result, it can be problematic if that operator is invoked on any other values, including uninitialized values.
However, there have been numerous reports of Thrust/CUB algorithms invoking user-defined operators on uninitialized, out-of-bound values or on repeatedly applying the reduction operator on items. See:
- https://github.com/NVIDIA/cccl/issues/780
- https://github.com/NVIDIA/cccl/issues/789
- https://github.com/NVIDIA/cccl/issues/822
- https://github.com/NVIDIA/cccl/issues/458
- https://github.com/NVIDIA/cccl/issues/801
There may be multiple root causes to this, particularly for the scan operation:
1. The CUDA algorithm implementations within CUB will have threads in a block do things like default construct objects, read past boundaries even if they don't ultimately participate in the work, or repeatedly apply the reduction operation more than once on some items. This is _fine_ for simple operators like `+` but for custom operators it can cause failures.
2. The decoupled look-back algorithm processes windows of 32 tile states at a time. It waits until all tile states of this window are not flagged as `SCAN_TILE_INVALID` anymore. I.e., each tile state of that window contains one of the following three values: (1) a tile's aggregate, (2) a tile's inclusive prefix, or (3) an out-of-bounds value, which is some "uninitialized" value. The decoupled look-back algorithm will then process this window, computing a `TailSegmentedReduce` over this window to compute the inclusive prefix. `TailSegmentedReduce` starts from the "rightmost" tile for which the inclusive prefix is available computing the reduction operation over this window. If there's no inclusive prefix amongst any of the tile states, it will just compute the aggregate across this window and continue processing the next window. We need to investigate whether `TailSegmentedReduce` is mathematically correct, or whether the reduction operation is applied on some invalid data (e.g., invoking the reduction operation on an aggregate and an item `x`, where `x` was already incorporated in the aggregate, hence breaking the invariant the user has on their data).
3. The decoupled look-back algorithm initializes the the first 32 tile states with the `SCAN_TILE_OOB` flag, to make sure we can process the full window that includes the very first "valid" tile state without having to account for reading beyond the beginning of tile states. These out-of-bounds tiles states are processed by `TailSegmentedReduce` and, hence, are passing invalid data to the reduction operation.
### Describe the solution you'd like
We should eliminate all cases where Thrust/CUB algorithms will invoke user-defined operators on out-of-bounds or uninitialized values.
### Tasks
- [ ] https://github.com/NVIDIA/cccl/issues/5017
- [ ] Make sure `thrust::set_intersection` invokes the reduction operator only with valid data
- [ ] Make sure `thrust::reduce_by_key`/`cub::DeviceReduce::ReduceByKey` invokes the reduction operator only with valid data
### Describe alternatives you've considered
_No response_
### Additional context
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.