NVIDIA / NVIDIA/cccl

[BUG]: `CUB::DeviceRunLengthEncode::Encode` is unclear about supported output types

Open
#400 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

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

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.