cuda核函数传入__device__函数指针的使用例子

cuda的global函数里面可以调用__device__函数，在有特殊需要的时候，还可以把__device__函数作为参数传入到一个__global__函数中
在cuda里面不能像c++那样简单地传入函数的指针，需要在传入前对函数的指针做一些包装。
例如

typedef double(*funcFormat)(int,char);

这里面double表示函数的返回值，int,char是函数的参数列表，所有满足这种格式的函数都可以用这种函数类型指代。
上面的funcFormat是一种函数类型，在这里可以把它理解成一种数据类型，从此funcFormat代表的就是double(*)(int,char)，后面声明__global__核函数的参数的时候，也直接使用furcFormat就行。
根据需要写好一个__device__函数之后，例如这样

__device__ double testFunc(int x,char y)
{
    //to do
}

这个函数testFunc仍然不能直接使用，需要用上面typedef出来的数据类型，把它声明成静态的函数指针

__device__ funcFormat staticTestFunc=testFunc;

，即使是这样，它仍然不能被直接使用，在传入函数指针的时候还需要把它转换成一个host函数指针，用cuda提供的cudaMemcpyFromSymbol接口来转换。

cudaMemcpyFromSymbol(&hostFunc,staticTestFunc,sizeof(funcFormat));

下面是一个具体的例子

#include<iostream>

//声明函数指针的类型
typedef double(*funcFormat)(int);

__device__ double testFunc(int x)
{
    return x*2+1;
}

//把函数指针声明成静态指针
__device__ funcFormat staticTestFunc=testFunc;

__global__ void cudaTest(double* gpuData,funcFormat func)
{
    gpuData[threadIdx.x]=(*func)(threadIdx.x);
}

int main()
{
    double* gpuData;
    cudaMalloc((void**)&gpuData,sizeof(double)*10);
    //新建一个位于host的同样的函数类型
    funcFormat hostFunc;
    //把之前指定过的静态指针复制到host部分
    cudaMemcpyFromSymbol(&hostFunc,staticTestFunc,sizeof(funcFormat));
    cudaTest<<<1,10>>>(gpuData,hostFunc);
    //把数据复制到cpu内存里面
    double cpuData[10];
    cudaMemcpy(cpuData,gpuData,sizeof(double)*10,cudaMemcpyDeviceToHost);
    for(int i=0;i<10;++i) std::cout<<cpuData[i]<<std::endl;
    return 0;
}