Search code examples
c++cudathrust

Error using Thrust Zip iterator with device functor


I'm working in a Linux environnement with Cuda and Thrust I got a standard code and i need to transform it Because of this i try to use use zip_iterator for the first time and i have some issues with it.

error: function "gridTransform::operator()" cannot be called with the given argument list

Does someone know how to resolve this issues . Thanks

thrust::host_vector<float> X_h;
thrust::host_vector<float> Y_h;
thrust::host_vector<float> Z_h;


typedef thrust::tuple<float,float,float>                           Float3;

typedef thrust::device_vector<float>::iterator                     FloatIterator;
typedef thrust::tuple<FloatIterator, FloatIterator, FloatIterator> FloatIteratorTuple;
typedef thrust::zip_iterator<FloatIteratorTuple>                  Float3Iterator;

struct gridTransform 
{
    float k;
    gridTransform(float _k)
    {
        k=_k;
    }
    __host__ __device__ Float3 operator()(Float3 &p1) const
    {
        float lx1=roundf(thrust::get<0>(p1)/k);
        float ly1=roundf(thrust::get<1>(p1)/k);
        float lz1=roundf(thrust::get<2>(p1)/k);

        return Float3(lx1,ly1,lz1);
    }
};

void Gpu_Functions::Decimation()
{
    thrust::device_vector<float> X_d(X_h);
    thrust::device_vector<float> Y_d(Y_h);
    thrust::device_vector<float> Z_d(Z_h);


    // Storage for result of each transformations
    thrust::device_vector<Float3> result(N);
    // Now we'll create some zip_iterators for A and B
    Float3Iterator A_first = thrust::make_zip_iterator(make_tuple(X_d.begin(), Y_d.begin(), Z_d.begin()));
    Float3Iterator A_last = thrust::make_zip_iterator(make_tuple(X_d.begin(), Y_d.begin(), Z_d.begin()));
   
    // vertices transformation generate duplicate vertices
    thrust::transform(A_first, A_last, result.begin(), gridTransform(0.01)); 

   // sort vertices to bring duplicates together
    thrust::sort(result.begin(), result.end());

    // find unique vertices and erase redundancies
   result.erase(thrust::unique(result.begin(), result.end()), result.end());
   
}

Solution

  • gridTransform::operator() needs to take a const Float3 &. I guess Thrust doesn't implement a cast from their internal thrust::detail::tuple_of_iterator_references<float &, float &, float &> to the non-const tuple.

    The reason for this design decision is probably that for a non-const reference the behavior would be somewhat unexpected: Writing something to this tuple would probably not result in a change to the input vectors in device memory, but only to some temporary tuple in registers.