在device_vector上使用thrust :: iterator_adaptor时出现问题

Question

我正在基于基础迭代器定义一个iterator_adaptor。整个过程在处理host_vectors时有效，但是当我将它应用于device_vectors时，编译器会抛出一个错误：对非const的引用的初始值必须是左值。

#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;
}

请不要通过指出使用推力提供的花式迭代器来实现相同效果的方法来回答。我需要开发一个迭代器来调整以预定义方式跨越晶格的类。这是给出错误的准系统代码。如果在主机中，代码将编译自定义的自适应迭代器，对于设备向量，代码将编译为设备向量的标准迭代器。当我在设备向量上使用自定义迭代器时，编译器会在上面示例中指示的位置抛出错误“对非const的引用的初始值必须是左值”。我用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;
...

而在代码中我现在向基础迭代器添加了一个模板参数，它接受了Container。所以现在我有

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

这解决了编译问题。为什么原始代码没有编译对我来说仍然是一个谜，但这提供了一个解决方案。