iterator for vector of structures in thrust

Question

I'm trying to get access to vector elements in this manner

struct point
{
    unsigned int x;
    unsigned int y;
};

...
thrust::device_vector<point> devPoints(hPoints.begin(), hPoints.end());

for(thrust::device_vector<point>::iterator iter = devPoints.begin(); iter != devPoints.end(); iter++) 
{
    std::cout << iter->x << " " << iter->y << " " << std::endl; (1)
}

device_vector was initialized properly. I get following errors:

error: expression must have pointer type (at 1)
error: no suitable user-defined conversion from "const thrust::detail::normal_iterator<thrust::device_ptr<point>>" to "thrust::device_ptr<point>" exists
          detected during instantiation of "Pointer thrust::experimental::iterator_facade<Derived, Pointer, Value, Space, Traversal, Reference, Difference>::operator->() const [with Derived=thrust::detail::normal_iterator<thrust::device_ptr<point>>, Pointer=thrust::device_ptr<point>, Value=point, Space=thrust::detail::cuda_device_space_tag, Traversal=thrust::random_access_traversal_tag, Reference=thrust::device_reference<point>, Difference=ptrdiff_t]"

What am I doing wrong?

Answer 1

Ok this one was a bit more complicated than I expected :)
Here are the results of my investigations:

Your problem comes from thrust's implementation. Thrust uses a type called device_reference which, as its documentation says: http://wiki.thrust.googlecode.com/hg/html/classthrust_1_1device__reference.html

device_reference acts as a reference-like object to an object stored in device memory. device_reference is not intended to be used directly; rather, this type is the result of deferencing a device_ptr. Similarly, taking the address of a device_reference yields a device_ptr.

However, there are some cases when we are dealing implicitly with device_reference. For example, when a device_reference is passed as a parameter to functions waiting for POD (more or less what you are trying to do with operator<<), the following problem appears:

Another common case where a device_reference cannot directly be used in place of its referent object occurs when passing them as parameters to functions like printf which have varargs parameters. Because varargs parameters must be Plain Old Data, a device_reference to a POD type requires a cast when passed to printf:

Having said that, all you have to do is to cast your device_reference to the POD you're handling. In your case, you'd do:

for(thrust::device_vector<point>::iterator iter = devPoints.begin(); iter != devPoints.end(); iter++)  {
   std::cout << (static_cast<point>(*iter)).x << " " << (static_cast<point>(*iter)).y << std::endl;
}

In my opinion, this is not the most elegant solution, I'd rather use the std::copy algorithm to print the content of your point class. Thus, I've written a small example file, using your point class and printing it using three different ways:

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <cstdlib>
#include <algorithm>
#include <iostream>

struct point
{
        unsigned int x;
        unsigned int y;
};

__host__
point getRandomPoint() {
        point p;

        p.x = rand();
        p.y = rand();

        return p;
}

__host__
std::ostream& operator<< (std::ostream& os, const point& p) {
        os << "[ " << p.x << " ; " << p.y << " ]";
        return os;
}

int main() {
        // fill the host_vector with random points
        thrust::host_vector<point> hPoints(512);
        thrust::generate(hPoints.begin(), hPoints.end(), getRandomPoint);

        // copy hPoints content to device memory
        thrust::device_vector<point> devPoints(hPoints.begin(), hPoints.end());

        // first way 
        for(thrust::device_vector<point>::iterator iter = devPoints.begin(); iter != devPoints.end(); iter++)  {
                std::cout << (static_cast<point>(*iter)).x << " " << (static_cast<point>(*iter)).y << std::endl;
        }

        // second way
        for(thrust::device_vector<point>::iterator iter = devPoints.begin(); iter != devPoints.end(); iter++)
        {
                std::cout << *iter << std::endl;
        }

        // third way
        std::copy(devPoints.begin(), devPoints.end(), std::ostream_iterator< point >(std::cout, "  $  ") );

        return 0;
}

Now, it's up to you to choose the one you prefer!

Answer 2

std::cout << iter->x << " " << iter->y << " " << std::endl; 
                                      ^^^^

:))