Compilation error with iterator_adaptor

[recke-a]

Hi,

I started this thread in the CUDA forum https://forums.developer.nvidia.com/t/issues-when-compiling-code-using-thrust-iterator-adaptor-with-nvcc/252669/4 and I was asked by Robert Crovella to file it here.

I tried to use a customized iterator (slightly changed code from the iterator_adaptor example) and I happened to get stuck with a compiler error about a deleted function.

The problem appeared with different settings now, but the first time after I updated to CUDA 12.1 including Thrust. I tried to boil it down to a simple example:

#include <iostream>
#include <thrust/iterator/iterator_adaptor.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/gather.h>
#include <thrust/execution_policy.h>

template<typename Iterator>
  class finite_sequence_repeat_iterator
    : public thrust::iterator_adaptor<
        finite_sequence_repeat_iterator<Iterator>, 
        Iterator
      >
{
  public:
    typedef thrust::iterator_adaptor<
      finite_sequence_repeat_iterator<Iterator>,
      Iterator
    > super_t;

    __host__ __device__
    finite_sequence_repeat_iterator(const Iterator &x, int n) : super_t(x), begin(x), n(n) {}

    friend class thrust::iterator_core_access;
  private:

    unsigned int n;

    const Iterator begin;

    __host__ __device__
    typename super_t::reference dereference() const
    {
      return *(begin + (this->base() - begin) % n);
    }
}; 

// function 

template <typename IteratorType>
finite_sequence_repeat_iterator<IteratorType> 
make_infinite_repeat_iterator(IteratorType start_iterator, unsigned int n)
{
    return finite_sequence_repeat_iterator<IteratorType>(start_iterator, n);
}

int main()
{
    thrust::device_vector<unsigned int> x(100, 0);
    thrust::device_vector<unsigned int> y(100, 0);

    thrust::sequence(thrust::device, x.begin(), x.end());

    std::cout << "\nVorher:\n>";
    for (auto it = x.begin(); it != x.end(); ++it) std::cout << (*it) << " ";
    std::cout << "<\n";

    auto ci = thrust::make_counting_iterator<unsigned int>(0);
    auto rep = make_infinite_repeat_iterator(ci, 5);

    thrust::gather(thrust::device, rep, rep+100, x.begin(), y.begin());

    std::cout << "\n>";
    for (auto it = y.begin(); it != y.end(); ++it) std::cout << (*it) << " ";
    std::cout << "<\n";
}

I used nvcc -arch=sm_52 --std=c++20 --extended-lambda --default-stream per-thread -c UnitTestRepeatIterator.cu -o UnitTestRepeatIterator to compile the program.

Which then led to the following error message:

/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_adaptor.h(206): error: function "finite_sequence_repeat_iterator<Iterator>::operator=(const finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>> &) [with Iterator=thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>]" (declared implicitly) cannot be referenced -- it is a deleted function
        m_iterator = static_cast<base_type>(m_iterator + n);
                   ^
          detected during:
            instantiation of "void thrust::iterator_adaptor<Derived, Base, Value, System, Traversal, Reference, Difference>::advance(thrust::iterator_adaptor<Derived, Base, Value, System, Traversal, Reference, Difference>::difference_type) [with Derived=thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>>>, Base=finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>>, Value=unsigned int, System=thrust::device_system_tag, Traversal=thrust::use_default, Reference=thrust::device_reference<unsigned int>, Difference=thrust::use_default]" at line 170 of /usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_facade.h
            instantiation of "void thrust::iterator_core_access::advance(Facade &, Facade::difference_type) [with Facade=thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>>>]" at line 379 of /usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_facade.h
            instantiation of "Derived &thrust::iterator_facade<Derived, Value, System, Traversal, Reference, Difference>::operator+=(thrust::iterator_facade<Derived, Value, System, Traversal, Reference, Difference>::difference_type) [with Derived=thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>>>, Value=unsigned int, System=thrust::device_system_tag, Traversal=thrust::random_access_traversal_tag, Reference=thrust::device_reference<unsigned int>, Difference=signed long]" at line 520 of /usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_facade.h
            instantiation of "Derived thrust::operator+(const thrust::iterator_facade<Derived, Value, System, Traversal, Reference, Difference> &, Derived::difference_type) [with Derived=thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>>>, Value=unsigned int, System=thrust::device_system_tag, Traversal=thrust::random_access_traversal_tag, Reference=thrust::device_reference<unsigned int>, Difference=signed long]" at line 327 of /usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/iterator/iterator_facade.h
            instantiation of "thrust::iterator_facade<Derived, Value, System, Traversal, Reference, Difference>::reference thrust::iterator_facade<Derived, Value, System, Traversal, Reference, Difference>::operator[](thrust::iterator_facade<Derived, Value, System, Traversal, Reference, Difference>::difference_type) const [with Derived=thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>>>, Value=unsigned int, System=thrust::device_system_tag, Traversal=thrust::random_access_traversal_tag, Reference=thrust::device_reference<unsigned int>, Difference=signed long]" at line 118 of /usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/system/cuda/detail/transform.h
            [ 9 instantiation contexts not shown ]
            instantiation of "OutputIt thrust::cuda_cub::transform_if(thrust::cuda_cub::execution_policy<Derived> &, InputIt, InputIt, StencilInputIt, OutputIt, TransformOp, Predicate) [with Derived=thrust::cuda_cub::par_t, InputIt=thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>>>, OutputIt=thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, StencilInputIt=thrust::cuda_cub::__transform::no_stencil_tag, TransformOp=thrust::cuda_cub::identity, Predicate=thrust::cuda_cub::__transform::always_true_predicate]" at line 334 of /usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/system/cuda/detail/transform.h
            instantiation of "OutputIt thrust::cuda_cub::transform_if(thrust::cuda_cub::execution_policy<Derived> &, InputIt, InputIt, OutputIt, TransformOp, Predicate) [with Derived=thrust::cuda_cub::par_t, InputIt=thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>>>, OutputIt=thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, TransformOp=thrust::cuda_cub::identity, Predicate=thrust::cuda_cub::__transform::always_true_predicate]" at line 353 of /usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/system/cuda/detail/transform.h
            instantiation of "OutputIt thrust::cuda_cub::transform(thrust::cuda_cub::execution_policy<Derived> &, InputIt, InputIt, OutputIt, TransformOp) [with Derived=thrust::cuda_cub::par_t, InputIt=thrust::permutation_iterator<thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>>>, OutputIt=thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, TransformOp=thrust::cuda_cub::identity]" at line 53 of /usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/system/cuda/detail/gather.h
            instantiation of "ResultIt thrust::cuda_cub::gather(thrust::cuda_cub::execution_policy<Derived> &, MapIt, MapIt, ItemsIt, ResultIt) [with Derived=thrust::cuda_cub::par_t, MapIt=finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>>, ItemsIt=thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, ResultIt=thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>]" at line 42 of /usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/detail/gather.inl
            instantiation of "OutputIterator thrust::gather(const thrust::detail::execution_policy_base<DerivedPolicy> &, InputIterator, InputIterator, RandomAccessIterator, OutputIterator) [with DerivedPolicy=thrust::cuda_cub::par_t, InputIterator=finite_sequence_repeat_iterator<thrust::counting_iterator<unsigned int, thrust::use_default, thrust::use_default, thrust::use_default>>, RandomAccessIterator=thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>, OutputIterator=thrust::detail::normal_iterator<thrust::device_ptr<unsigned int>>]" at line 67 of UnitTestRepeatIterator.cu

1 error detected in the compilation of "UnitTestRepeatIterator.cu".

Interestingly, when I use thrust::copy with this iterator, no error message appears. It seems that it has something to do with a missing or deleted copy assignment operator.

In fact, the only difference to the example code for thrust::iterator_adaptor is that the return *(begin + (this->base() - begin) / n); is now return *(begin + (this->base() - begin) % n); in line 38.

Of course, I can find a workaround by using a mixture of transform_iterator and permutation_iterator to achieve the same effect, but iterator_adaptor is meant to work, too, which is the reason why I filed this issue.

I hope someone has an idea of how to solve this.

Kind regards Andreas

[jrhemstad]

Hey @recke-a. Luckily (or perhaps unluckily) I've run into this kind of problem numerous times and so I knew what to look for 🙂.

The problem is the const Iterator begin member of your finite_sequence_repeat_iterator.

When a class has a const data member, the compiler implicitly deletes the copy assignment operator: example. This has caused me so much headache in the past that I almost never use const data members in classes anymore.

Several Thrust algorithms use copy-assignment (many in places where they probably shouldn't), so when your iterator has a deleted copy assignment iterator, it will fail.

The good news is that the fix is easy. Your example compiles just fine if I remove the const: https://godbolt.org/z/P84axYjTc

[recke-a]

Dear jrhemstad,

many thanks for your helpful response. In fact, having a const member variable begin here actually does not make any sense. It is private and the only functions that access it are the constructor and the dereference function which is defined as a const function and cannot change member variables.

I will send that to the CUDA guys, too to close the thread.

Kind regards Andreas