transform_iterator fails with callables whose arguments are passed by non-const l-value ref

[jrhemstad]

Minimal reproducer: https://godbolt.org/z/s4oPhqx3W

This fails because it cannot deduce the return type of the callable.

A workaround is to explicitly define the result_type in the callable: https://godbolt.org/z/MTM8b1zW9

However, I believe a better solution is to fix the result type deduction.

It currently fails because it is ultimately attempting to do:

using result_type = std::invoke_result_t<ref, int>;

This fails due to using int as the argument type instead of int&.

This stems from using thrust::iterator_value on the adapted iterator here: https://github.com/NVIDIA/thrust/blob/d50708f1be248f4a42e81fdce360cf368990582e/thrust/iterator/detail/transform_iterator.inl#L39-L42

Instead of getting the value_type of the iterator, I believe it should use iterator_reference:

    typedef typename thrust::detail::ia_dflt_help<
      Reference,
      thrust::detail::result_of_adaptable_function<UnaryFunc(typename raw_reference<typename thrust::iterator_reference<Iterator>::type>>::type)>
    >::type reference;

Note the use of raw_reference to avoid issues with device_vector::iterator::reference.

Example: https://godbolt.org/z/8eePb5Kvc

[alliepiper]

Can you share a bit more about the usecase here? Using transform_iterator with a functor that modifies the input sequence seems problematic, and I'm not sure this should be allowed.

[jrhemstad]

Can you share a bit more about the usecase here? Using transform_iterator with a functor that modifies the input sequence seems problematic, and I'm not sure this should be allowed.

I've been experimenting with creating a new fancy iterator that allows accessing a data member of an object. There have been several cases in the past where I've wanted to compute the result of something like a scan directly into into existing objects without having to materialize the intermediate scan result and then copy it into the existing objects.

One way I was trying to do this was via a transform iterator that takes an l-value ref of the object and returns an l-value ref of the member:

https://godbolt.org/z/jd1obbbWG

Generally I'd agree with you that taking an l-value ref in a transform function is a bad idea as it almost certainly means side-effects. In this case there aren't any side-effects. It's moreso to create a different kind of "transform_output_iterator".

[bernhardmgruber]

I have ran into this issue from multiple directions by now and may propose a solution at some point in the future.

[bernhardmgruber]

For a transform_iterator with a projecting transformation, like:

struct foo {
  int x, y;
};

struct access_x {
  _CCCL_HOST_DEVICE int& operator()(foo& f) const noexcept {
    return f.x;
  }
};

auto it = thrust::make_transform_iterator(base_it, access_x{});
auto& ref = *it;

there are four scenarios to consider:

iterator_system of base_it is	base_it::reference is a	implementation	*it called on	comment	projection can work?
host	l-value reference	can pass reference to transformation	host		yes
host	l-value reference	can pass reference to transformation	device	segfault or cross-system access at runtime	no
device	wrapped reference	can unwrap reference and pass to transformation	host	segfault or cross-system access at runtime	no
device	wrapped reference	can unwrap reference and pass to transformation	device		yes

The problem is that whether the implementation is correct depends on where the iterator is dereferenced from. We could have divergent code paths for host and device code for the dereference implementation, but this would mean we call the transformation function with different types in host and device code. This in turn could lead to the transformation function to return different types as well (e.g. cuda::std::identity) or fail to compile in either host or device code. However, thrust::transform_iterator<>::reference needs to be the same type for host and device code, so this type needs to match whatever the transformation returns.

[bernhardmgruber]

For iterators which are wrapped pointers, this

auto it = thrust::make_transform_iterator(thrust::raw_pointer_cast(base_it), access_x{});

almost works. It just loses the device system property of base_it: if base_it was a device_ptr then its iterator system was device. But the raw_pointer_cast strips this information and the transform_iterator suddenly has a host iterator system.

This can be worked around by explicitly specifying the iterator system (requires exposing a System template parameter on transform_iterator):

auto it = thrust::transform_iterator<access_x, foo*, int&, int, thrust::device_system_tag>(base_it, access_x{});

[bernhardmgruber]

For the purpose of projections, counting_iterator<T*> can almost solve this: #4014. Adding a strided_iterator<T*, stride> on top, should do it.

[bernhardmgruber]

cuda::transform_iterator will solve this.

[bernhardmgruber]

The following code works now:

#include <thrust/iterator/counting_iterator.h>
#include <cuda/iterator>
#include <vector>
#include <algorithm>

struct foo{
  int v;
  int w;
};

int main(){
  std::vector<foo> f(10);

  auto it = cuda::make_transform_iterator(f.begin(), &foo::v);
  std::copy(thrust::make_counting_iterator(0), 
            thrust::make_counting_iterator(10), 
            it);

  assert(f[9].v == 9);
  assert(f[9].w == 0);

}

https://godbolt.org/z/YcqGa8bz9