[FEA] Add iterator wrapping a T* with a custom element stride

[bernhardmgruber]

To run CUB/thrust algorithms on a single data member of an array of structs, we would need an iterator that wraps a pointer to that data member (T*), but advances a different stride than sizeof(T).

Contrived example use case:

struct rgb { float r, g, b; };

std::vector<rgb> image = ...;
float* first_red = &image[0].r;
// ++first_red would advance to image[0].g
auto it = striding_iterator<float, sizeof(rgb)>{first_red};
// ++it would jump to image[1].r
auto sum_red = thrust::reduce(host, it, it + image.size());

[miscco]

I am wondering whether it would be better to ensure that thrust::transform_iterator can work with projections like ranges::transform

[bernhardmgruber]

I am wondering whether it would be better to ensure that thrust::transform_iterator can work with projections like ranges::transform

I agree, but I tried that in the past (#2006) and the problem is that thrust::transform_iterator wraps a lot of thrust::device_pointer<T> returning thrust::device_reference<T>, so we would need to make cuda::std::invoke(ptr, &rgb::x) work on a thrust::device_pointer<T>.

[bernhardmgruber]

It turns out counting_iterator<Incrementable> can almost solve this, since the Incrementable can be a pointer and #4014 proposes a custom step. We just need to add a dereference twice.