Unify turning iterators into pointers

[bernhardmgruber]

cuda::std::contiguous_iterator<thrust::device_vector<int>::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.

[bernhardmgruber]

@miscco you may have some ideas on how to achieve this. For example, I wonder whether ::cuda::std::__can_to_address is actually enough in all our use cases. thrust::unwrap_contiguous_iterator can be entirely replaced by cuda::std::to_address.

[bernhardmgruber]

@miscco why do we need ::cuda::std::__can_to_address? Are their cuda::std::contiguous_iterators where ::cuda::std::to_address does not work?