NVIDIA / NVIDIA/cccl

Unify turning iterators into pointers

Open
#9,893 2 comments 0 reactions 2 assignees Claimed by @bernhardmgruber View on GitHub
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

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.