Cuda的函数指针

Question

我试图让这样somtehing（其实我需要写一些集成功能）的CUDA

我试过，但它并没有工作 - 它不仅造成。

Error: Function pointers and function template parameters are not supported in sm_1x.

float f1(float x) { 
    return x; 
} 

__global__ void tabulate(float lower, float upper, float p_function(float), float*result){ 
    for (lower; lower < upper; lower++) { 
       *result = *result + p_function(lower); 
     } 
} 

int main(){ 
     float res; 
    float* dev_res; 

     cudaMalloc((void**)&dev_res, sizeof(float)) ; 

    tabulate<<<1,1>>>(0.0, 5.0, f1, dev_res); 
    cudaMemcpy(&res, dev_res, sizeof(float), cudaMemcpyDeviceToHost) ; 

    printf("%f\n", res); 
    /************************************************************************/ 
    scanf("%s"); 

    return 0; 

} 

Answer 1

为了摆脱编译错误，编译代码时必须使用-gencode arch=compute_20,code=sm_20作为编译器参数。但你很可能有一些运行时的问题：

从CUDA编程指南采取http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#functions

Function pointers to __global__ functions are supported in host code, but not in device code. Function pointers to __device__ functions are only supported in device code compiled for devices of compute capability 2.x and higher.

It is not allowed to take the address of a __device__ function in host code.

这样你就可以有这样的事情（改编自“FunctionPointers”样品）：

//your function pointer type - returns unsigned char, takes parameters of type unsigned char and float 
typedef unsigned char(*pointFunction_t)(unsigned char, float); 

//some device function to be pointed to 
__device__ unsigned char 
Threshold(unsigned char in, float thresh) 
{ 
    ... 
} 

//pComputeThreshold is a device-side function pointer to your __device__ function 
__device__ pointFunction_t pComputeThreshold = Threshold; 
//the host-side function pointer to your __device__ function 
pointFunction_t h_pointFunction; 

//in host code: copy the function pointers to their host equivalent 
cudaMemcpyFromSymbol(&h_pointFunction, pComputeThreshold, sizeof(pointFunction_t)) 

然后，您可以将h_pointFunction作为参数传递给您的内核，该内核可以使用它来调用您的__device__函数。

//your kernel taking your __device__ function pointer as a parameter 
__global__ void kernel(pointFunction_t pPointOperation) 
{ 
    unsigned char tmp; 
    ... 
    tmp = (*pPointOperation)(tmp, 150.0) 
    ... 
} 

//invoke the kernel in host code, passing in your host-side __device__ function pointer 
kernel<<<...>>>(h_pointFunction); 

希望这是有道理的。总之，它看起来像你将不得不将f1函数改为__device__函数并遵循类似的过程（typedefs不是必需的，但它们确实使代码更好），以将其作为有效的函数指针主机端传递给你的内核。我还建议给FunctionPointers CUDA示例看看

Answer 2

即使你可以编译这段代码（见@Robert Crovella的回答）这个代码将无法正常工作。由于主机编译器无法确定函数地址，因此无法从主机代码传递函数指针。