NVIDIA / NVIDIA/cccl

[FEA]: Allow BlockRunLengthDecode to be initialized with striped offsets and items

Open
#3,844 1 comment 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 request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Area

CUB

### Is your feature request related to a problem? Please describe.

Currently `cub:: BlockRunLengthDecode::InitWithRunOffsets()` fills shared memory with offsets and items in a blocked fashion which potentially causes bank conflicts:

https://github.com/NVIDIA/cccl/blob/06b5541683390e97d0ec84dc1ccbe60fa762f652/cub/cub/block/block_run_length_decode.cuh#L311-L317

I understand that a blocked layout is the default for CUB block algorithms and that in this case the assumption is that `RUNS_PER_THREAD` is small as each run can produce a lot of data and bigger values would use significant amounts of shared memory. Also this initialization is needed for the constructor taking run lengths which have to be scanned first as the scan output is blocked as well.

But given that the offsets potentially come from global memory (and aren't necessarily aligned for vectorized access considering load balancing of block-workloads), this makes the user pay twice (global load/exchange and shared store) for the blocked layout that is not necessary.

### Describe the solution you'd like

To avoid having a constructor or algorithm or similar for every possible layout (striped, warp-striped, ...), I propose a generalized constructor that takes not only the offsets and items, but also the ranks of the offsets and items in a third array/span so any data layout can be used.

### Describe alternatives you've considered

While I don't expect a significant performance uplift for typical workloads (given a small `RUNS_PER_THREAD` even a `cub::LoadDirectBlocked()` for the global load should be fine due to L1 caching and the bank conflicts should not result in significant serialization either), the ease of mind when using this constructor and the ease of implementing it made me write this issue either way.

### Additional context

_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.