Проблема при использовании thrust :: iterator_adaptor на device_vector

Question

Я работаю над определением iterator_adaptor на основе базового итератора.Все это работает при работе с host_vectors, но когда я применяю его к device_vectors, компилятор выдает ошибку: начальное значение ссылки на non-const должно быть 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;
}

Пожалуйста, не отвечайте, указав способы достижения того же эффекта, используя причудливые итераторы, которые обеспечивает тяга.Мне нужно разработать итератор, который адаптирует класс, который переходит через решетку заранее определенным образом.Это тот самый код, который выдает ошибку.Код компилируется для специально адаптированного итератора, если он находится в хосте, и для стандартного итератора векторов устройств для векторов устройств.Когда я использую пользовательский итератор для векторов устройств, компилятор выдает ошибку «начальное значение ссылки на неконстантное значение должно быть lvalue» в месте, указанном в примере выше.Я компилирую его с помощью nvcc main.cu.nvcc версия 9.0, gcc версия 7.3.0, версия 1.9.0

Answer 1

Я думаю, что нашел решение.Проблема связана с определением типов элементов для базового итератора.Замена приведенного выше кода на этот

#include <thrust/iterator/iterator_adaptor.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
template <typename T>
struct BoxIterator
{

  unsigned int m_loc;
  typedef typename decltype(T().begin())::difference_type difference_type;
  typedef typename decltype(T().begin())::pointer pointer;
  typedef typename decltype(T().begin())::reference reference;
  typedef typename decltype(T().begin())::value_type value_type;
  typedef typename decltype(T().begin())::iterator_category 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<T> &other) const
  {
    return m_loc == other.m_loc;
  }

  __host__ __device__ bool equal(const BoxIterator<T> &other) const
  {
    return m_loc == other.m_loc;
  }

  __host__ __device__
      difference_type
      distance_to(const BoxIterator<T> &other) const
  {
    return other.m_loc - this->m_loc;
  }
  __host__ __device__
      BoxIterator<T>
      operator+(int n) const
  {
    BoxIterator<T> tmp = *this;
    tmp.m_loc += n;
    return tmp;
  }
  __host__ __device__
      BoxIterator<T>
      operator-(int n) const
  {
    BoxIterator<T> tmp = *this;
    tmp.m_loc -= n;
    return tmp;
  }
  __host__ __device__
      BoxIterator
      operator-(const BoxIterator<T> other) const
  {
    BoxIterator<T> tmp = *this;
    tmp.m_loc -= other.m_loc;
    return tmp;
  }
  __host__ __device__
  BoxIterator(const BoxIterator<T> &other)
  {
    m_loc = other.m_loc;
  }
  __host__ __device__
      BoxIterator &
      operator=(const BoxIterator<T> &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.
  }

  __host__ __device__
      typename super_t::difference_type
      distance_to(const SubVolumeIterator<LatticeIt, Container> &other) const
  {
    return (other.base().m_loc - this->base().m_loc); //+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<thrust::host_vector<double>> 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;
    }
  }

  {
    BoxIterator<thrust::device_vector<double>> bit;
    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;
}

приводит к коду, который компилируется и работает правильно.Основное различие заключается в том, как типы элементов определяются в базовом BoxIterator.Первоначально я определил их как

typedef double value_type; 
typedef double& reference;
...

, тогда как в коде теперь я добавил аргумент шаблона в базовый итератор, который принимает контейнер.Так что теперь у меня есть

typedef typename decltype(T().begin())::value_type value_type;
...

Это исправляет проблему компиляции.Почему исходный код не скомпилирован, для меня все еще остается загадкой, но это предлагает решение.