Using polymorphic functors inside functions in Thrust

Question

I need a function that computes some "math" function of several variables on GPU. I decided to use Thrust and its zip_iterator to pack variables in a tuple, and implement my math function as a functor foк for_each. But I'd like to have a universal function that can compute different "math" functions. So, I need to pass this functor in the function.

As I think, to do this task, I should implement some simple hierarchy (with the only base class) of functors with different versions of operator()(Tuple t). For example, functors can be like these:

struct Func {
    template <typename Tuple>
    __host__ __device__
    void operator()(Tuple arg) {
        thrust::get<1>(arg) = 0;
    }
};

strust f1 : public Func {
   ...
};

strust f2 : public Func {
   ...
};

The question is how can I properly pass required functor into the function, and how to define such a function? This function can be like:

thrust::host_vector evaluateFunction(Func func, thrust::host_vector point) {
    thrust::host_vector result(point.size());

    thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(point.begin(), result.begin())),
    thrust::make_zip_iterator(thrust::make_tuple(point.end(), result.end())),
                 func);
    return result;
}

but with such a definition I cannot pass f1 or f2 into it. How can I define it properly?

Answer 1

It appears that you want to pass a user-defined or user-selectable function to thrust, for evaluation in a functor. Based on that, I think a possible answer is very similar to what is posted here, however I'll make some other comments.

1. I don't think you were actually asking about this, but I think polymorphism expressed in host classes, objects of which that are eventually passed to the device for use in a thrust::device_vector would be essentially impossible, because It is not allowed to pass as an argument to a __global__ function an object of a class with virtual functions. Although it's not entirely obvious, this is in effect what I would be doing if I created an appropriate array of objects on the host and then attempted to use them polymorphically in a thrust::device_vector. At least, I could not fight my way through the jungle. You can use polymorphism in device code, of course, assuming the objects are correctly created on the device so that their virtual function tables can be properly assigned. This question and comments may also be of interest. If you want to see an example of using polymorphism on device objects, take a look at my answer here. It demonstrates the idea of using the object itself to define the function you wish to perform on it. (Although, again, I don't think that is what you had in mind.)
2. Based on your edited question, it appears that you'd like to pass the address of a function, for use in a functor. This is also (at least in a straightforward manner) impossible with CUDA, because it is not allowed to take the address of a __device__ function in host code. This limitation can be worked around, after a fashion, by calling an "initialization kernel" which fills in a table of device function pointers for use in device code, but I think the net effect is no better than what I propose in 3 below.
3. You can pass a user-selectable index of a function, for use in device code.

Here's an example which demonstrates this last idea:

#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/copy.h>
#include <thrust/for_each.h>
#include <thrust/iterator/zip_iterator.h>
#include <iostream>
#include <math.h>


__host__ __device__ float f1(float x)
{
  return sinf(x);
}

__host__ __device__ float f2(float x)
{
  return cosf(x);
}



struct select_functor
{
  unsigned fn;

  select_functor(unsigned _fn) : fn(_fn) {};
  template <typename Tuple>
  __host__ __device__
  void operator()(const Tuple &t) {
    if (fn == 1) thrust::get<1>(t)  = f1(thrust::get<0>(t));
    else if (fn == 2) thrust::get<1>(t)  = f2(thrust::get<0>(t));
    else thrust::get<1>(t) = 0;
  }
};


int main(void)
{
  unsigned ufn = 1;
  const unsigned N = 8;
  thrust::device_vector<float> data(N), result(N);
  // initilaize to some values
  thrust::sequence(data.begin(), data.end(),  0.0f, (float)(6.283/(float)N));
  std::cout<< "x: " << std::endl;
  thrust::copy(data.begin(), data.end(), std::ostream_iterator<float>(std::cout, " "));
  thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(data.begin(), result.begin())),thrust::make_zip_iterator(thrust::make_tuple(data.end(),result.end())),select_functor(ufn));
  std::cout<< std::endl << "sin(x): " << std::endl;
  thrust::copy(result.begin(), result.end(), std::ostream_iterator<float>(std::cout, " "));
  ufn = 2;
  thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(data.begin(), result.begin())),thrust::make_zip_iterator(thrust::make_tuple(data.end(),result.end())),select_functor(ufn));
  std::cout<< std::endl << "cos(x): " << std::endl;
  thrust::copy(result.begin(), result.end(), std::ostream_iterator<float>(std::cout, " "));
  std::cout<< std::endl;
  return 0;
}