I am working on defining an iterator_adaptor based on a base iterator. The whole thing works when working on host_vectors, but when I apply it to device_vectors, the compiler throws an error: initial value of reference to non-const must be an lvalue.
#include <thrust/iterator/iterator_adaptor.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
struct BoxIterator{
unsigned int m_loc;
typedef int difference_type;
typedef double* pointer;
typedef double& reference;
typedef double value_type;
typedef thrust::random_access_device_iterator_tag iterator_category;
__host__ __device__
BoxIterator() : m_loc(0){}
__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;
}
__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;
}
};
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;
__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); //+this->base().m_loc); // this gives an error: initial value of reference to non-const must be an lvalue
// when thrust::copy is used on a device_vector. Compiles fine with a host_vector.
}
};
int main()
{
thrust::host_vector<double> HV(100);
thrust::device_vector<double> DV(100);
thrust::device_vector<double> DV1(100);
BoxIterator bit;
SubVolumeIterator<decltype(bit), decltype(HV)> HIt(bit, HV, 1);
SubVolumeIterator<decltype(bit), decltype(HV)> HIt_end(bit + 20, HV, 1);
thrust::fill(HIt, HIt_end, 5.); // this compiles fine
for (int i = 1; i < 21; ++i)
{
std::cout << HV[i] << std::endl;
}
{
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
}
{
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
thrust::copy(DV.begin()+1,DV.begin()+21, HV.begin()+1);
for (int i = 1; i < 21; ++i)
{
std::cout << HV[i] << std::endl;
}
}
return 0;
}
Please do not answer by pointing out ways to achieve the same effect using the fancy iterators that thrust provides. I need to develop an iterator that adapts a class that steps over a lattice in a predefined way. This is the barebone code that gives the error. The code compiles for a custom adapted iterator if in host, and for the standard iterator of device vectors for device vectors. When I use the custom iterator on device vectors the compiler throws the error "initial value of reference to non-const must be an lvalue" at the location indicated in the example above. I compile it with nvcc main.cu. nvcc version 9.0, gcc version 7.3.0, thrust version 1.9.0
double
) in both locations (the example code basically suggests to use defaults for parameters after the 2nd), and remove the__device__
decorator from the class constructor (its not necessary for your usage), and the code compiles correctly and runs without error for me. Furthermore, thefill
results appear correct betweenDV
andHV
at the end. – Robert Crovella Apr 9 '19 at 1:11.cu
file and not a.cpp
file – talonmies Apr 9 '19 at 5:34