I am trying to use thrust::iterator_adaptor to create a specialized iterator that “walks” over an array defined over a lattice in a specialized way. Everything works out when applied to arrays that are defined in CPU (host_memory). However, when I use it on arrays defined in GPU (device_memory), the code does not compile. Instead, I get an “error: initial value of reference to non-const must be an lvalue” when compiling the function that dereferences the pointer. I created a simplified version of the code that still exhibits the problem. The comment lines in the code go over the details of the problem. Please note that the problem is not with the algorithm, and the code below is a very stripped down version, just to illustrate the problem.
#include <thrust/iterator/iterator_adaptor.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
struct BoxIterator{ // this is basically a counting iterator.
// It implements the basic functions (++, advance,...) expected from an iterator
// It is the simplest code that still triggers the exception.
// data member
unsigned int m_loc;
// type members
typedef int difference_type;
typedef double* pointer;
typedef double& reference;
typedef double value_type;
typedef thrust::random_access_device_iterator_tag iterator_category;
// basic constructor
__host__ __device__
BoxIterator() : m_loc(0){}
//copy and assignment constructors
__host__ __device__
BoxIterator(const BoxIterator & other)
{
m_loc = other.m_loc;
}
__host__ __device__
BoxIterator & operator=(const BoxIterator & other)
{
m_loc = other.m_loc;
return *this;
}
// bare minimum members
__host__ __device__
void operator++()
{
++m_loc;
}
__host__ __device__
void advance(int n)
{
m_loc += n;
}
__host__ __device__
void operator--()
{
advance(-1);
}
__host__ __device__
void operator+=(int n)
{
advance(n);
}
__host__ __device__
void begin()
{
m_loc = 0;
}
__host__ __device__
bool operator==(const BoxIterator & other) const
{return m_loc==other.m_loc;}
__host__ __device__
bool equal(const BoxIterator & other) const
{
return m_loc==other.m_loc;
}
__host__ __device__
difference_type distance_to(const BoxIterator & other) const
{
return other.m_loc - this->m_loc;
}
__host__ __device__
BoxIterator operator+(int n)
{
BoxIterator tmp = *this;
tmp.m_loc += n;
return tmp;
}
}; // end of BoxIterator
// this is the adapted iterator.
template <typename LatticeIt, typename Container>
class SubVolumeIterator : public thrust::iterator_adaptor<SubVolumeIterator<LatticeIt, Container>,
LatticeIt
>
{
public:
typedef thrust::iterator_adaptor<SubVolumeIterator<LatticeIt, Container>,
LatticeIt
>
super_t;
// Basically, the idea is to use the iterator provided by the Container, and dereference it on a position given by the
// LatticeIt.
__host__
SubVolumeIterator(const LatticeIt &It, Container &FAB, int N) : super_t(It),
v(FAB.begin()),
offset(N) {}
friend class thrust::iterator_core_access;
private:
decltype(Container().begin()) v;
int offset;
__host__ __device__
typename super_t::reference
dereference() const
{
return *(v + offset); // In reality, we are interested in dereferencing v on a location
// that depends on this->base()
// but here for simplicity, we just dereference it
// on a constant offset.
// When LatticeIt is the standard iterator,
// it compiles and runs for
// Containers in host and device (provided line 16
// is modified to read random_access_host_....
// When LatticeIt is our own (e.g, BoxIterator as defined above),
// it compiles and runs on host Containers,
// but it does not compile on device Containers,
// giving instead the error " initial value of
// reference to non-const must be an lvalue"
}
}; // end of SubVolumeIterator
int main()
{
thrust::device_vector<double> DV(100);
BoxIterator bit;
{
SubVolumeIterator<decltype(DV.begin()), decltype(DV)> DIt(DV.begin(), DV, 5);
SubVolumeIterator<decltype(DV.begin()), decltype(DV)> DIt_end(DV.begin() + 20, DV, 5);
thrust::fill(DIt,DIt_end , -5.); // this compiles fine, with LatticeIt
// being the standard iterator over the device_vector
}
{ // if the code in this block is commented out, it compiles fine.
SubVolumeIterator<decltype(bit), decltype(DV)> DIt(bit, DV, 5);
SubVolumeIterator<decltype(bit), decltype(DV)> DIt_end(bit + 20, DV, 5);
thrust::fill(DIt,DIt_end , -5.); // this throws the error
}
return 0;
}
I am using nvcc 9.0, thrust version is 1.9.0, and gcc version is 7.3.0. My gut feeling is that there is something missing in BoxIterator that messes up code introspection, but have no idea what that may be. Note that if instead we use host_vectors, and of course
change line 16 so that the iterator_category is random_access_host_iterator_tag, everything compiles without issues.