c++ - 在 Thrust 的函数内部使用多态仿函数

Question

我需要一个函数来计算 GPU 上几个变量的一些“数学”函数。我决定使用 Thrust 及其 zip_iterator 将变量打包到一个元组中，并将我的数学函数实现为仿函数 foк for_each。但我想要一个可以计算不同“数学”函数的通用函数。所以，我需要在函数中传递这个仿函数。

我认为，要完成这项任务，我应该使用不同版本的 operator()(Tuple t) 实现一些简单的仿函数层次结构(只有基类)。例如，仿函数可以是这样的:

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

strust f1 : public Func {
   ...
};

strust f2 : public Func {
   ...
};

问题是如何正确地将所需的仿函数传递给函数，以及如何定义这样的函数？这个函数可以是这样的:

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;
}

但是对于这样的定义，我无法将 f1 或 f2 传递给它。我怎样才能正确定义它？

Answer 1

您似乎想将用户定义或用户可选择的函数传递给 thrust，以便在仿函数中进行计算。基于此，我认为可能的答案与发布的内容非常相似 here ，但是我会提出一些其他意见。

1. 我不认为你真的在问这个，但我认为 polymorphism在主机类中表示，最终传递给设备以在 thrust::device_vector 中使用的对象基本上是不可能的，因为 It is not allowed to pass as an argument to a __global__ function an object of a class with virtual functions.虽然不是很明显，但如果我在主机上创建一个适当的对象数组然后尝试在 thrust::device_vector 中多态地使用它们，这实际上就是我要做的事情。至少，我无法在丛林中杀出一条血路。当然，您可以在设备代码中使用多态性，前提是对象已在设备上正确创建，以便可以正确分配它们的虚函数表。 This question and comments也可能感兴趣。如果您想查看在设备对象上使用多态性的示例，请查看我的回答 here .它演示了使用对象本身来定义您希望对其执行的功能的想法。 (不过，我认为这不是您的初衷。)
2. 根据您编辑的问题，您似乎想要传递一个函数的地址，以便在仿函数中使用。这对于 CUDA 也是不可能的(至少以直接的方式)，因为 it is not allowed to take the address of a __device__ function in host code.这个限制可以通过调用一个“初始化内核”来绕过，它填充了一个用于设备代码的设备函数指针表，但我认为净效果并不比我在下面的 3 中提出的更好。
3. 您可以传递一个用户可选择的函数索引，以供在设备代码中使用。

这是一个演示最后一个想法的例子:

#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;
}