NVIDIA / NVIDIA/cccl

[BUG]: GPU scan uses wrong aggregation type

Open
#5,668 0 comments 0 reactions 0 assignees 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

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

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.