Use `cub::DeviceTransform::Transform` in `thrust::tabulate`

[bernhardmgruber]

First, we need to

• Estimate effort of supporting thrust::tabulate by cub::DeviceTransform::Transform #5918

[bernhardmgruber]

So, this is a bit harder. We can express tabulate as a transformation of a counting iterator range to the output range. The counting iterator is not contiguous though, so we fall back to the prefetch algorithm. However, we want to use the vectorized storing from the vectorized algorithm.

One way we can go forward is to allow non-contiguous iterators as input to the vectorized algorithm, where we would just load the data scalar by scalar. We may also want to combine the vectorized and prefetching algorithm, since the vectorized one already includes the prefetching one as a runtime fallback. However, the prefetching algorithm has the advantage of being runtime tuneable to achieve certain bytes in flight and thus a runtime number of elements per thread. Combining this may be difficult and time consuming.

[miscco]

Could we detect, whether the return value of iter is a reference and if not take the fast path anyway?

[bernhardmgruber]

Could we detect, whether the return value of iter is a reference and if not take the fast path anyway?

The check I chose is now whether the iterator is contiguous or not. I only regard contiguous iterators in inputs for whether we can vectorize loads or not. If the iterator is not contiguous, I just load it serially into a local array. Then I process serially and store vectors again if possible.