Issues when compiling code using thrust::iterator_adaptor with nvcc

[Andreas_Lubeck]

Hi there,

I stumbled over a weird problem since I updated to CUDA 12 and the latest Thrust release.
It appears it has something to do with hidden iterator semantics with iterator_adaptor, and with the settings for C++ 2017 standard.

I actually took the example for iterator_adaptor to build an iterator that actually repeats the first n elements like 1,2,3,…,n,1,2,3,…,n, 1,2,3,…n.

// derive repeat_iterator from iterator_adaptor
template
class infinite_repeat_iterator
: public thrust::iterator_adaptor<
infinite_repeat_iterator, // the first template parameter is the name of the iterator we’re creating
Iterator // the second template parameter is the name of the iterator we’re adapting
// we can use the default for the additional template parameters
>
{
public:
// shorthand for the name of the iterator_adaptor we’re deriving from
typedef thrust::iterator_adaptor<
infinite_repeat_iterator,
Iterator
> super_t;
host device
infinite_repeat_iterator(const Iterator &x, int n) : super_t(x), begin(x), n(n) {}
// befriend thrust::iterator_core_access to allow it access to the private interface below
friend class thrust::iterator_core_access;
private:
// repeat each element of the adapted range n times
unsigned int n;
// used to keep track of where we began
const Iterator begin;
// it is private because only thrust::iterator_core_access needs access to it
host device
typename super_t::reference dereference() const
{
return *(begin + (this->base() - begin) % n);
}
};

The compiler expands this and says this …
/include/thrust/detail/tuple.inl(479): error: function “custom_iterators::infinite_repeat_iterator::operator=(const custom_iterators::infinite_repeat_iterator<thrust::detail::normal_iterator<thrust::device_ptr>> &) [with Iterator=thrust::detail::normal_iterator<thrust::device_ptr>]” (declared implicitly) cannot be referenced – it is a deleted function
cons& operator=(const cons& u) { head = u.head; return *this; }
^
It seems that the implicit assignment operator has been deleted for some reason. And that I need to add something to my customized iterator definition.
I am stuck. I hope someone could help here. It must be something simple, but I cannot find it.

Many thanks in advance
Andreas

[Robert_Crovella]

when posting code here, please use the formatting tools available. As an example: click the pencil icon under your post to edit the post. Select the code. Press the </> at the top of the edit window, then save your changes.

I also suggest providing a complete example that someone could compile. I also suggest providing a link to the example you started with. Finally, if this is for a version of thrust that does not ship with the CUDA toolkit you are using, you may wish to file a thrust issue instead.

[Andreas_Lubeck]

Dear Robert, I am sorry for the mistake. Here comes the code boiled down to the problem.

File UnitTestRepeatIterator.cu

#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_53 --std=c++17 --extended-lambda --default-stream per-thread -c UnitTestRepeatIterator.cu -o UnitTestRepeatIterator
to compile the program.

Which then led to the following error (truncated):

/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);

I have an NVIDIA GeForce GTX 970 card with Driver Version: 530.30.02 and CUDA Version: 12.1. I know that the Maxwell architecture is marked as deprecated from CUDA version 11. Maybe I have to install an older version, but the compiler error message appears to be connected to something else.

Actually, I can use thrust::copy to copy the contents of the repeat iterator into a device_vector without complaints by the compiler. But when I use it for the map argument in thrust::gather as above, the compiler reports an error. In both cases, the repeat iterator is used according to the InputIterator model.
The assignment operator for the iterator is the problem here, but I cannot figure out why that is.

Thank you in advance
Andreas

[Robert_Crovella]

I guess you are working from this.

Mainly what I see is you have changed / to % here:

@Andreas_Lubeck

return *(begin + (this->base() - begin) % n);

I don’t see a problem with your code. And as you say, this problem pops up when you use your iterator with gather but not with some other operations. So I think it might be a thrust issue, and you might want to file one. I also note the statement from the iterator_adapter page I linked:

However, the large set of iterator semantics which must be satisfied for algorithm compatibility can make iterator_adaptor difficult to use correctly. Unless you require the full expressivity of iterator_adaptor , consider building a custom iterator through composition of existing higher-level fancy iterators instead.

Perhaps someone else will be able to help.

It’s not related to your issue, but sm_53 is not the right choice for a GTX 970. You should be using sm_52.

[Andreas_Lubeck]

@Andreas_Lubeck

Actually, I can use thrust::copy to copy the contents of

Dear Robert,

I filed this issue now at Thrust’s github site as you suggested: https://github.com/NVIDIA/thrust/issues/1937

Best regards
Andreas

[Andreas_Lubeck]

Dear Robert,

we found the solution.
It is very simple: the const member variable begin is the problem.

 const Iterator begin;

The copy-assignment operator is deleted, but is needed by thrust::gather and not by thrust::copy.

Actually, the example for thrust::iterator_adaptor is misleading. Anyway, just removing the const in that definition is the simple magic that was needed here.

Kind regards
Andreas