How to advance iterator in thrust function

Question

I'm doing some study on thrust. But I didn't understand how to get the value of an iterator point to.

An example code is like:

#include <thrust/for_each.h>
#include <thrust/device_vector.h>
#include <iostream>
#include <vector>
using namespace std;

class ADD
{
private:
    typedef typename thrust::device_vector<int>::iterator PTR;
public:
    ADD(){}
    ~ADD(){}
    void setPtr(PTR &ptr)
    {this->ptr=ptr;}
    __host__ __device__
    void operator()(int &x)
    {
            // note that using printf in a __device__ function requires
            // code compiled for a GPU with compute capability 2.0 or
            // higher (nvcc --arch=sm_20)
            x+=add();
    }
    __host__ __device__
    int add()
    {return *ptr++;}
private:
    PTR ptr;
};
int main()
{
    thrust::device_vector<int> d_vec(3);
    d_vec[0] = 0; d_vec[1] = 1; d_vec[2] = 2;
    thrust::device_vector<int>::iterator itr=d_vec.begin();
    ADD *addtest=new ADD();
    addtest->setPtr(itr);
    thrust::for_each(d_vec.begin(), d_vec.end(), *addtest);
    for(int i=0;i<3;i++)
            cout<<d_vec[i]<<endl;
    return 0;
}

When I compile this using nvcc -arch=sm_20 test.cu I got the following warning:

test.cu(28): warning: calling a host function("thrust::experimental::iterator_facade<thrust::detail::normal_iterator<thrust::device_ptr<int> > , thrust::device_ptr<int> , int, thrust::detail::cuda_device_space_tag, thrust::random_access_traversal_tag, thrust::device_reference<int> , long> ::operator *") from a __device__/__global__ function("printf_functor::add") is not allowed

test.cu(28): warning: calling a host function("thrust::experimental::iterator_facade<thrust::detail::normal_iterator<thrust::device_ptr<int> > , thrust::device_ptr<int> , int, thrust::detail::cuda_device_space_tag, thrust::random_access_traversal_tag, thrust::device_reference<int> , long> ::operator *") from a __device__/__global__ function("printf_functor::add") is not allowed

I cannot get this to compile. How can I solve this problem?

Answer 1

@Gang.Wang: I think you just mixing up 2 different things: all STL-like functionality including for_each, device_vector iterators etc. is just a "facade" which exists on the host only.

While operator() contains the actual GPU code which is compiled to CUDA kernel and applied to each element of your vector in parallel. Hence, device_vector::iterators are not accessible from your functor.