我正在开发一个需要将任意函数传递给CUDA核心的CUDA应用程序。由于为每种可能情况声明函数指针并将它们传递到内核中会很麻烦(超过50个不同的函数),而且它们都是像
我想出了以下代码,这是一种可以在 NVCC 上编译并运行良好的最小实现。然而,通过这种方式,我只能实现具有1个变量的函数。是否有办法扩展我的代码来处理类似于
提前感谢!
sin(x)/y
这样的基本函数组合,因此我希望 CUDA 核心具有一些最小的 Lambda-表达式功能。由于设备代码尚不支持 C++11 特性(就我所知),而且我在网上找不到任何相关信息,所以我决定自己教授表达式模板,并实现一些简单的 lambda-表达式规则以传递到内核中。我想出了以下代码,这是一种可以在 NVCC 上编译并运行良好的最小实现。然而,通过这种方式,我只能实现具有1个变量的函数。是否有办法扩展我的代码来处理类似于
sin(_x) + _y
的函数组合?提前感谢!
#include<math.h>
#ifdef __CUDACC__
#define HOST_DEVICE __host__ __device__
#else
#define HOST_DEVICE
#endif
struct Id {};
template <typename Op, typename Left, typename Right>
struct BinaryOp
{
Left left;
Right right;
HOST_DEVICE BinaryOp(Left t1, Right t2) : left(t1), right(t2) {}
HOST_DEVICE double operator() (double x) {
return Op::apply(left(x), right(x));
}
};
template <typename Op, typename Arg>
struct UnaryOp
{
Arg arg;
HOST_DEVICE UnaryOp(Arg t1) : arg(t1) {}
HOST_DEVICE double operator() (double x) {
return Op::apply(arg(x));
}
};
template <>
struct UnaryOp<Id, double>
{
HOST_DEVICE UnaryOp() {}
HOST_DEVICE double operator() (double x) {
return x;
}
};
struct Sin
{
HOST_DEVICE static double apply(double x) {
return sin(x);
}
};
struct Plus
{
HOST_DEVICE static double apply(double a, double b) {
return a + b;
}
};
template <typename Left, typename Right>
BinaryOp<Plus, Left, Right> operator+ (Left lhs, Right rhs) {
return BinaryOp<Plus, Left, Right>(lhs, rhs);
}
template <typename Arg>
UnaryOp<Sin, Arg> _sin(Arg arg) {
return UnaryOp<Sin, Arg>(arg);
}
template <class T>
__global__ void test(T func, double x) {
printf("%e\n", func(x));
}
int main ()
{
UnaryOp<Id, double> _x;
double x = 1.0;
test<<<1, 1>>>(_sin(_x) + _x, x);
cudaDeviceSynchronize(); // Needed or the host will return before kernel is finished
return 0;
}