Unify turning iterators into pointers
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
`cuda::std::contiguous_iterator::iterator` is `false`:
Thrust vector iterators are not `cuda::std::contiguous_iterator` according to the C++20 concept. They only fulfill `thrust::is_contiguous_iterator_v` (which can be extended by a user via `THRUST_PROCLAIM_CONTIGUOUS_ITERATOR`). This is a problem, because several places in CCCL guard a fast path by `cuda::std::contiguous_iterator`, for example the proposed `BlockPrefetch` (#9723), but also load vectorization in `BlockLoad::Load`. These optimizations are not taken for thrust vectors, since they fail the trait.
Repro: https://cuda.godbolt.org/z/e6d3Kh1Yb
I believe we should come up with a guideline on how to use `cuda::std::contiguous_iterator`, `thrust::is_contiguous_iterator_v`, `THRUST_PROCLAIM_CONTIGUOUS_ITERATOR`, `THRUST_NS_QUALIFIER::unwrap_contiguous_iterator`, `::cuda::std::__can_to_address` in our code base, based on the semantics that we can turn an iterator into a pointer for fast paths, and how a user can enroll their iterators into this mechanism.
We should then deprecate all the Thrust machinery. We need to keep `THRUST_PROCLAIM_CONTIGUOUS_ITERATOR` in user code working for a while.
Contributor guide
Assessment
This issue has not been assessed yet.