[BUG]: Cub's ReduceByKey passing unexpected values to reduction_op
- 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
I was trying to utilize the [cub::DeviceReduce::ReduceByKey](https://nvidia.github.io/cccl/cub/api/structcub_1_1DeviceReduce.html#_CPPv4I0000000EN3cub12DeviceReduce11ReduceByKeyE11cudaError_tPvR6size_t18KeysInputIteratorT21UniqueOutputIteratorT20ValuesInputIteratorT25AggregatesOutputIteratorT22NumRunsOutputIteratorT12ReductionOpT9NumItemsT12cudaStream_t) function and noticed that my implementation was resulting in `cudaGetLastError` returning an illegal memory access error. Specifically, in my program, the values passed into the `Reduce` represent indexes, and the custom reduction function would use those passed in indexes to read values from a buffer. However, I noticed that for some threads in the last scheduled thread block, this function would pass in garbage values into this reduction function, causing it to access indexes greater than the buffer's size, resulting in illegal memory accesses.
### How to Reproduce
You can find an example code recreating this behaviour at: [https://godbolt.org/z/1eE96EYbo](https://godbolt.org/z/1eE96EYbo). When I compile and run this code using the command: `nvcc cub_bug.cu -o cub_bug && ./cub_bug`, I get the output that:
```
Running experiment with num records of 1048576
First 5 Key records: 0 1 2 3 4
First 5 Value records: 0 0 0 0 0
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [43,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [44,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [45,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [46,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [47,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [48,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [49,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [50,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [51,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [52,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [53,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [54,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [55,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [56,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [57,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [58,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [59,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [60,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [61,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [62,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [63,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [96,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [97,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [98,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [99,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [100,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [101,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [102,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [103,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [104,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [105,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [106,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [107,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [108,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [109,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [110,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [111,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [112,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [113,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [114,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [115,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [116,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [117,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [118,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [119,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [120,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [121,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [122,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [123,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [124,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [125,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [126,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [127,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [64,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [65,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [66,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [67,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [68,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [69,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [70,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [71,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [72,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [73,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [74,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [75,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [76,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [77,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [78,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [79,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [80,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [81,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [82,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [83,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [84,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [85,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [86,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [87,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [88,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [89,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [90,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [91,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [92,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [93,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [94,0,0] Assertion `lower_value < num_records` failed.
cub_bug.cu:51: unsigned long CustomReductionOperator::operator()(const unsigned long &, const unsigned long &) const: block: [1365,0,0], thread: [95,0,0] Assertion `lower_value < num_records` failed.
CUDA error: device-side assert triggered
```
Note that I am running this code on a AWS g6.xlarge instance with an Ubuntu 22.04 image and an L4 GPU. I have included the result of the `nvidia-smi` command with the cuda and driver versions.
### Expected behavior
The CustomReductionOperator should never have an invalid index/value passed to it. For now, I have just added a check in my code to ignore any invalid values and I see the right result being produced but ideally, I shouldn't need to have a special case like this.
### Reproduction link
https://godbolt.org/z/1eE96EYbo
### Operating System
Ubuntu Linux 22.04
### nvidia-smi output
```
NVIDIA-SMI 570.86.10 Driver Version: 570.86.10 CUDA Version: 12.8 |
|-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA L4 Off | 00000000:31:00.0 Off | 0 |
| N/A 29C P8 15W / 72W | 1MiB / 23034MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+
+-----------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=========================================================================================|
| No running processes found |
```
### NVCC version
```
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Wed_Jan_15_19:20:09_PST_2025
Cuda compilation tools, release 12.8, V12.8.61
Build cuda_12.8.r12.8/compiler.35404655_0
```
Contributor guide
Assessment
This issue has not been assessed yet.