[FEA]: Add a CUB BlockLoad using prefetching
- 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
Assessment
This issue has not been assessed yet.