如何将设备函数作为输入参数传递给主机端函数？答案

【问题标题】：How to pass device function as an input argument to host-side function?如何将设备函数作为输入参数传递给主机端函数？
【发布时间】：2013-08-28 10:57:21
【问题描述】：

我只想将设备函数作为主机函数的参数传递，当然，主机函数可以使用这个设备端函数启动一些内核。

我尝试了通常的 C++ 方式（通过指针/引用传递），CUDA 调试器告诉我内核无法启动。

更新：

我想做的是：

__host__ void hostfunction(int a, int (*DeviceFunction)(int))
{
   /...do something.../
   somekernel<<<blocks, threads>>>(int * in, DeviceFunction);
}

并使用以下命令启动主机：

hostfunction(x, &SomeDeviceFunctionTemplate<int>);

【问题讨论】：

你的问题并不完全清楚，至少对我来说是这样。如果您可以发布一个您尝试过的示例，其中内核未启动，那可能会有所帮助。
由于__host__函数不能获取__device__函数的地址，你基本上需要编写一个简短的__global__函数来获取感兴趣的__device__函数的地址，并且然后将其存储到内存中。然后，您的 __host__ 函数可以从内存中读取该函数指针，然后将其传递给 somekernel。

标签： pointers reference cuda nvidia

【解决方案1】：

这个例子可能很有趣：

$ cat t237.cu
#include <stdio.h>


__device__ int f1(){ printf("dev f1\n"); return 0;}
__device__ int f2(){ printf("dev f2\n"); return 0;}
__device__ int f3(){ printf("dev f3\n"); return 0;}

__device__ int *fptrf1 = (int *)f1;
__device__ int *fptrf2 = (int *)f2;
__device__ int *fptrf3 = (int *)f3;


__global__ void mykernel(int (*fptr)()){

  fptr();
  printf("executed\n");
}

int main(){

  int *hf1, *hf2, *hf3;
  cudaMemcpyFromSymbol(&hf1, fptrf1, sizeof(int *));
  cudaMemcpyFromSymbol(&hf2, fptrf2, sizeof(int *));
  cudaMemcpyFromSymbol(&hf3, fptrf3, sizeof(int *));
  mykernel<<<1,1>>>((int (*)())hf1);
  cudaDeviceSynchronize();
  mykernel<<<1,1>>>((int (*)())hf2);
  cudaDeviceSynchronize();
  mykernel<<<1,1>>>((int (*)())hf3);
  cudaDeviceSynchronize();
  return 0;
}
$ nvcc -arch=sm_20 -O3 -o t237 t237.cu
$ ./t237
dev f1
executed
dev f2
executed
dev f3
executed
[bob@cluster1 misc]$

我认为这与 Jared 的建议大致相符。正如他所提到的，这在主机代码中是不可能的：

&SomeDeviceFunctionTemplate<int>

假设SomeDeviceFunctionTemplate 指的是__device__ 函数。

【讨论】：

我能否以某种方式将其调整为模板代码，而不是您在main() 中的代码？还是没有办法为我计划使用的每个函数显式定义设备端指针？
一些链接here的示例使用模板。
因此，this example 将全局指针变量的需要替换为全局设置内核的需要，它将设备函数的指针复制到启动器选择的位置。这没有多大帮助......继续阅读。关于内核地址的问题是可用的。但是我该如何利用它来获取任意设备函数的地址呢？

【解决方案2】：

如果您能发布一个您正在尝试执行的操作的示例，将会很有帮助，但要检查的一件事是您正在 Fermi (sm_20) 或更高版本上编译和运行，因为较旧的 GPU 不支持非内联函数调用。

检查您设备的计算能力（需要 2.0 或更高版本）并检查您的 nvcc 命令行（需要 -arch=sm_20 或更高版本，或 -gencode 等效版本）。

【讨论】：