NVIDIA / NVIDIA/cccl

[FEA]: BlockLoad/BlockStore to/from shared memory

Open
#885 8 comments 0 reactions 1 assignee Assigned to @pauleonix View on GitHub
cub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

For some algorithms it makes no sense to have a whole block of data in registers at once. For others a local buffer is bad due to dynamic indexing. For these it would be great to have versions of `cub::BlockLoad` and `cub::BlockStore` that skip the step of loading the values from shared memory into registers and just provide them in shared memory s.t. each thread of a warp has all its values on the same bank.

One idea would be to create versions of the `Load` / `Store` methods that don't take a local array and then have a device function or method to access data in `TempStorage`, s.t. the user doesn't have to care about the explicit memory layout and implementation details can be changed without breaking user code.

Naturally all this would only work for algorithms that store all data in shared memory at the same time. But as the core implementation is already existing, the implementation should hopefully be straightforward?

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.