Writing a simple thrust functor operating on some zipped arrays

Question

I am trying to perform a thrust::reduce_by_key using zip and permutation iterators. i.e. doing this on a zipped array of several 'virtual' permuted arrays. I am having trouble in writing the syntax for the functor density_update.

But first the setup of the problem.

Here is my function call:

thrust::reduce_by_key(  dflagt,
                        dflagtend,
                        thrust::make_zip_iterator( 
                            thrust::make_tuple(
                                thrust::make_permutation_iterator(dmasst, dmapt), 
                                thrust::make_permutation_iterator(dvelt, dmapt),
                                thrust::make_permutation_iterator(dmasst, dflagt),
                                thrust::make_permutation_iterator(dvelt, dflagt)
                            )
                        ),
                        thrust::make_discard_iterator(),
                        danswert,
                        thrust::equal_to<int>(),
                        density_update()
                  ) 

dmapt, dflagt are of type thrust::device_ptr<int> and dvelt , dmasst and danst are of type thrust::device_ptr<double>.

(They are thrust wrappers to my raw cuda arrays)

The arrays mapt and flagt are both index vectors from which I need to perform a gather operation from the arrays dmasst and dvelt.

After the reduction step I intend to write my data to the danswert array. Since multiple arrays are being used in the reduction, obviously I am using zip iterators.

My problem lies in writing the functor density_update which is binary operation.

struct density_update
{
  typedef thrust::device_ptr<double> ElementIterator;
  typedef thrust::device_ptr<int>   IndexIterator;
  typedef thrust::permutation_iterator<ElementIterator,IndexIterator> PIt;

  typedef thrust::tuple< PIt , PIt , PIt, PIt> Tuple;
  __host__ __device__
  double operator()(const Tuple& x , const Tuple& y)
  {   
      return    thrust::get<0>(*x) * (thrust::get<1>(*x) - thrust::get<3>(*x)) + \
                thrust::get<0>(*y) * (thrust::get<1>(*y) - thrust::get<3>(*y));
  }
};

The value being returned is a double . Why the binary operation looks like the above functor is not important. I just want to know how I would go about correcting the above syntactically. As shown above the code is throwing a number of compilation errors. I am not sure where I have gone wrong.

I am using CUDA 4.0 on GTX 570 on Ubuntu 10.10

Answer 1

density_update should not receive tuples of iterators as parameters -- it needs tuples of the iterators' references.

In principle you could write density_update::operator() in terms of the particular reference type of the various iterators, but it's simpler to have the compiler infer the type of the parameters:

struct density_update
{
  template<typename Tuple>
  __host__ __device__
  double operator()(const Tuple& x, const Tuple& y)
  {   
    return thrust::get<0>(x) * (thrust::get<1>(x) - thrust::get<3>(x)) + \
           thrust::get<0>(y) * (thrust::get<1>(y) - thrust::get<3>(y));
  }
};