Iterators should be iterators #3767
Description
During the discussion of a CUB PR, we started wondering what we can require of an iterator in host code. An argument was made that some iterators are only advancable in device code, so we have to work around this. I object. An iterator must have a type that satisfies the requirements of an iterator, which even for the most basic iterator includes incrementability.
We further wondered, whether we can rely on being able to compare iterators or compute their difference. As an example, we discussed an iterator applying an offset stored in device memory before dereferencing it:
template <typename It>
struct offset_iterator {
int* offset; // somewhere in device memory
It base;
__host__ __device__ auto operator+=(int difference) {
base += difference;
}
__device__ auto& dereference() {
return *(base + *offset);
}
friend bool operator==()(offset_iterator a, offset_iterator b) {
return a.base + *a.offset == b.base + *b.offset
}
...
};Such an iterator can be advanced in host and device code, since we can delegate to advancing the base iterator, but we cannot compare it without touching device memory.
Does such a type qualify as iterator in host code? Should it fulfill the corresponding iterator concepts? Should CUB support (and workaround) such types, which do not provide certain iterator APIs in host (or device) code?