[BUG]: Dereferencing a zip_iterator and converting to its value type leads to a compilation error

[andrewcorrigan]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Compile-time Error

Component

Thrust

Describe the bug

Casting the result of dereferencing a zip_iterator to its value type leads a compilation error for certain types, but not others. A minimal reproducer is shown below

How to Reproduce

Please see the godbolt reproducer link below.

Defining T in the minimal reproducer to contain thrust::tuple<double> triggers a compilation error, while various other cases do not. The failures and non-failures are consistent between g++ and clang++

#define THRUST_DEVICE_SYSTEM 4
#include <thrust/tuple.h>
#include <thrust/device_vector.h>
#include <thrust/iterator/zip_iterator.h>

int main(int argc, char** argv)
{
    // FAILURES:
    using T = thrust::tuple<double>;
    //using T = thrust::tuple<thrust::tuple<double>,double>;

    // WORKS:
    //using T = double; 
    //using T = thrust::tuple<double,double>;
    //using T = thrust::tuple<thrust::tuple<double,double>,double>;

    thrust::device_vector<T> data(10);

    auto a = thrust::make_zip_iterator(data.begin());

    static_assert(std::is_same_v<thrust::tuple<T>, ::cuda::std::iter_value_t<decltype(a)>>);    // success

    auto b = a[0];

    static_assert(std::is_same_v<thrust::detail::tuple_of_iterator_references<thrust::device_reference<T>>, decltype(b)>);    // success

    auto c = thrust::tuple<T>(b);       // errors for the cases above

    return 0;
}

Expected behavior

It should compile, instead it errors if the the type T contains thrust::tuple<double>

In file included from /opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/detail/internal_functional.h:23,
                 from /opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/system/detail/generic/copy.inl:15,
                 from /opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/system/detail/generic/copy.h:30,
                 from /opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/detail/copy.inl:20,
                 from /opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/detail/copy.h:57,
                 from /opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/detail/allocator/copy_construct_range.h:17,
                 from /opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/detail/contiguous_storage.inl:16,
                 from /opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/detail/contiguous_storage.h:214,
                 from /opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/detail/vector_base.h:21,
                 from /opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/device_vector.h:20,
                 from <source>:3:
/opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/iterator/detail/tuple_of_iterator_references.h: In instantiation of 'U thrust::detail::maybe_unwrap_nested<U, T>::operator()(const T&) const [with U = double; T = thrust::device_reference<cuda::std::__4::tuple<double> >]':
/opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/iterator/detail/tuple_of_iterator_references.h:118:42:   required from 'constexpr cuda::std::__4::tuple<_Args1 ...> thrust::detail::tuple_of_iterator_references<Ts>::__to_tuple(cuda::std::__4::__tuple_indices<_Indices ...>) const [with Us = {double}; long unsigned int ...Id = {0}; Ts = {thrust::device_reference<cuda::std::__4::tuple<double> >}]'
  118 |     return {maybe_unwrap_nested<Us, Ts>{}(::cuda::std::get<Id>(*this))...};
      |             ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/iterator/detail/tuple_of_iterator_references.h:104:29:   required from here
  104 |     return __to_tuple<Us...>(typename ::cuda::std::__make_tuple_indices<sizeof...(Ts)>::type{});
      |            ~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/opt/compiler-explorer/libs/cccl/trunk/thrust/thrust/iterator/detail/tuple_of_iterator_references.h:36:12: error: cannot convert 'const thrust::device_reference<cuda::std::__4::tuple<double> >' to 'double' in return
   36 |     return t;
      |            ^
Compiler returned: 1

Reproduction link

https://godbolt.org/z/nhhf6xPra

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[miscco]

Thanks for the bug report,

This is a bit of a pickle, because tuple_of_iterator_references is never meant to be a type that is stored anywhere, but only used for conversions.

The code in question works fine if changed to auto c = thrust::tuple<T>(cuda::std::move(b));

Will investigate whether there is a simple workaround

[miscco]

So I looked into this a bit and I am not sure, whether this is something we actually want spend a lot of time on.

The underlying problem is that there are different conversion sequences at play here, which get the compiler confused whether it wants a tuple<tuple<A>> or just a tuple<A>

Again, tuple_of_iterator_reference is not meant to be used as a vocabulary type but is a cludge to help with device_reference

Is there a specific pattern this comes up with?

Note that again assigning directly works fine thrust::tuple<T> b = a[0];

[andrewcorrigan]

Thank you for looking into this and providing both an explanation and workarounds.

Is there a specific pattern this comes up with?

Yes. I have a helper function that converts the result of an iterator dereferencing to its value type, something like:

auto value(const auto& t)
{
    // `remove_reference_value_t` is defined recursively to filter out `device_reference` and convert `tuple_of_iterator_references` to `tuple`.
    return remove_reference_value_t<decltype(t)>(t);
}

and code that calls it like:

      auto t = value(t_start[i]);

Based on the workarounds you provided, I refactored it to something like the following.

      auto t = value(t_start, i);

That worked fine. I have thousands of calls like this with lots of different combinations of iterators so this is comprehensive.

Even with the workaround and the understanding that tuple_of_iterator_references is not a public type to be stored or manipulated, it still comes up a lot in basic user code and this was surprising that this broke depending on the underlying type. Also, this was working as recently as Thrust 1.11.0 (finally getting around to upgrading now...) and for as long as I can remember before that. I already spent a day looking into this and recognize the complexity of the conversion sequences you referenced, but I might try to see if I can find a fix to the underlying issue.

[miscco]

Did you try to move towards cuda::zip_iterator or cuda::zip_transform_iterator?

[andrewcorrigan]

I had not tried that. I wasn't even aware there was an alternative version in libcudacxx. But, thank you, I just tried it out and it handled all the cases correctly for the reproducer I sent. I will try it out for the entirety of my code as well.

[andrewcorrigan]

I was testing it further in my code and hit some compilation errors using OpenMP or TBB. It looks like cuda::zip_iterator has issues as well. I extracted a minimal reproducer. It gets a bit intricate but here are two representative failures for cuda::zip_iterator. For one of them thrust::zip_iterator actually works fine.

A case where both cuda::zip_iterator and thrust::zip_iterator fail

When copying from a doubly-nested zip_iterator, and using either the TBB or OpenMP device system, compilation fails. It works fine using the CPP device system.

https://godbolt.org/z/fn9z9xWP3

#define THRUST_DEVICE_SYSTEM 2 // OpenMP is broken
//#define THRUST_DEVICE_SYSTEM 3 // TBB is broken
//#define THRUST_DEVICE_SYSTEM 4 // CPP works

#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/copy.h>

#if 0 // ::cuda::std

#include <cuda/iterator>
#include <cuda/std/tuple>
template<typename... T> using my_tuple = ::cuda::std::tuple<T...>;
auto my_make_zip_iterator(auto... a) { return ::cuda::make_zip_iterator(a...); }

#else // ::thrust

#include <thrust/iterator/zip_iterator.h>
#include <thrust/tuple.h>
template<typename... T> using my_tuple = ::thrust::tuple<T...>;
auto my_make_zip_iterator(auto... a) { return thrust::make_zip_iterator(a...); }

#endif

int main(int argc, char** argv)
{
    using T = int;
    thrust::device_vector<T> a(10, 0);

    thrust::device_vector<my_tuple<my_tuple<T>>> b(a.size());

    thrust::copy_n(my_make_zip_iterator(my_make_zip_iterator(a.begin())),
                   a.size(),
                   b.begin());

    return 0;
}

A case where both cuda::zip_iterator fails but thrust::zip_iterator works

If the doubly-nested zip_iterator is 2x2, then it works with thrust:: but not with cuda:: using TBB or OpenMP. Both work using CPP

https://godbolt.org/z/n7GejvTnj

#define THRUST_DEVICE_SYSTEM 2 // OpenMP
//#define THRUST_DEVICE_SYSTEM 3 // TBB
//#define THRUST_DEVICE_SYSTEM 4 // CPP

#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/copy.h>

#if 1 // ::cuda::std -- FAILS

#include <cuda/iterator>
#include <cuda/std/tuple>
template<typename... T> using my_tuple = ::cuda::std::tuple<T...>;
auto my_make_zip_iterator(auto... a) { return ::cuda::make_zip_iterator(a...); }

#else // ::thrust -- WORKS

#include <thrust/iterator/zip_iterator.h>
#include <thrust/tuple.h>
template<typename... T> using my_tuple = ::thrust::tuple<T...>;
auto my_make_zip_iterator(auto... a) { return thrust::make_zip_iterator(a...); }

#endif

int main(int argc, char** argv)
{
    using T = int;
    thrust::device_vector<T> a(10, 0);

    thrust::device_vector<my_tuple<my_tuple<T,T>,my_tuple<T,T>>> b(a.size());

    thrust::copy_n(my_make_zip_iterator(my_make_zip_iterator(a.begin(), a.begin()),
                                        my_make_zip_iterator(a.begin(), a.begin()) ),
                   a.size(),
                   b.begin());

    return 0;
}

[andrewcorrigan]

I added this failure to this issue, since from the errors I'm seeing, it looks like it's the same underlying conversion error, but please let me know if this should be in a separate issue.