[BUG]: `CUB::DeviceRunLengthEncode::Encode` is unclear about supported output types
- 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
Something else
### Component
CUB
### Describe the bug
Currently we have no restrictions on the type of `UniqueOutputIteratorT` passed to `CUB::DeviceRunLengthEncode::Encode`
However, the implementation currently implicitly requires that both `InputIteratorT` and `UniqueOutputIteratorT` have the same value type due to the way tile predecessors are handled currently:
```cpp
// Load tile predecessor key in first thread
KeyOutputT tile_predecessor;
if (threadIdx.x == 0)
{
// if (tile_idx == 0)
// first tile gets repeat of first item (thus first item will not
// be flagged as a head)
// else
// Subsequent tiles get last key from previous tile
tile_predecessor = (tile_idx == 0) ? keys[0] : d_keys_in[tile_offset - 1];
}
```
Here, `key[0]` has type `KeyOutputT` while `d_keys_in` is of type `KeysInputIteratorT`
We should either enhance the implementation to support different output types or clarify the documentation.
### How to Reproduce
Use `CUB::DeviceRunLengthEncode::Encode` with a different output type
### Expected behavior
We should be clear whether this is allowed and implement it if desired
### Reproduction link
_No response_
### Operating System
_No response_
### nvidia-smi output
_No response_
### NVCC version
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.