[BUG]: GPU scan uses wrong aggregation type
- 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
Silent Failure
### Component
Thrust
### Describe the bug
I post this as a Thrust bug as I found it there, but the underlying bug is in `cub::BlockScan`. In the multiple-elements-per-thread case `cub::BlockScan` will use `cub::ThreadReduce` but not specify its `AccumT`. By default this type is chosen as whatever is returned by the binary operator applied to the values of the input type. This is inconsistent with `BlockScan` using the input/output/value type `T` for accumulation everywhere else.
Example member function of `BlockScan` (not the one used by Thrust in the reproducer, but they all have this problem):
https://github.com/NVIDIA/cccl/blob/d52ffc7393a35abf15d06c6fd6a96c681a0d186d/cub/cub/block/block_scan.cuh#L1062-L1074
Signature of `ThreadReduce`:
https://github.com/NVIDIA/cccl/blob/d52ffc7393a35abf15d06c6fd6a96c681a0d186d/cub/cub/thread/thread_reduce.cuh#L169-L173
(`cub::detail::ThreadScanExclusive` was not yet modernized to determine its accumulator type the same way as `cub::ThreadReduce`, but the new `cub::detail::ThreadScanExclusivePartial` was)
### How to Reproduce
0. See godbolt link below for exact reproducer
1. Define a custom binary operator with differing input and output types that has a deleted overload taking its output type as the first argument:
```c++
struct custom_op {
__host__ __device__ op_t operator()(value_t left, value_t right) {
return cuda::std::plus{}(left, right);
}
// This one should never be called!
__host__ __device__ op_t operator()(op_t, value_t) = delete;
};
```
2. Try to compile a device scan using this as the binary scan operator on an input range of `value_t`.
One can also create a runtime check by not deleting the overload and instead letting it return some constant. This then shows that this overload is actually used at runtime to compute valid outputs when one chooses a big enough input range that more than one thread works on valid output elements.
### Expected behavior
The given code should compile.
### Reproduction link
Compile time: https://godbolt.org/z/x4bd7M7zb
Run time: https://godbolt.org/z/4ohh3hYzn
### Operating System
_No response_
### nvidia-smi output
_No response_
### NVCC version
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.