NVIDIA / NVIDIA/cccl

[FEA]: Add a CUB BlockLoad using prefetching

Open
#3,126 1 comment 2 reactions 1 assignee Assigned to @gonidelis View on GitHub
cub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

Given the success we had with the prefetching implementation in `cub::DeviceTransform`, we should try to refactor this into one or several new `cub::BlockLoad` and validate its impact in at least one CUB algorithm.

- [x] Decide which block load algorithms would benefit from prefetching. I assume `BLOCK_LOAD_DIRECT` and possibly `BLOCK_LOAD_VECTORIZE`.
- [x] Add the equivalent of `cub::LoadDirectBlocked` etc. (3 overloads) using prefetching. Look at `cub::DeviceTransform` for how to prefetch to L2
- [x] Add a new enumerator to `cub::BlockLoadAlgorithm` and implement it in `cub::BlockLoad`
- [x] Examine benefits of using L1 prefetching (available only in CTK higher than 13.4)

**With all the above decided during the design process, this will now be used as a tracking issue for prefetching. See sub-issues below.**

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.