如何将函数作为 cuda 内核参数传递？答案

【问题标题】：How to pass a function as a cuda kernel parameter?如何将函数作为 cuda 内核参数传递？
【发布时间】：2019-12-07 04:17:19
【问题描述】：

我想创建一个通用 cuda 内核，它接受一个可调用对象作为参数（如 lambda 或函数）并调用它。

我无法将设备函数作为参数传递给 cuda 内核。

我有计算能力 3.5 的 cuda 9.2。我在 Debian 10 上使用 gcc 9.3。

我试过这个，用nvcc -arch=sm_35 --expt-extended-lambda main.cu -o test编译：

    __host__ __device__ void say_hello()
    {
        printf("Hello World from function!\n");
    }

    template<class Function>
    __global__ void generic_kernel(Function f)
    {
        f();
    }

    int main() 
    {
            // this is working
        generic_kernel<<<1,1>>>([]__device__(){printf("Hello World from lambda!\n");});
        cudaDeviceSynchronize();

            // this is not working!
        generic_kernel<<<1,1>>>(say_hello); 
        cudaDeviceSynchronize();

        return 0;
    }

我希望同时看到 Hello World from function! 和 Hello World from lambda!，但我只看到来自 lambda 的消息。

【问题讨论】：

第二个例子是将主机函数传递给内核，这显然是错误的，还有其他问题
我认为say_hello 是一个设备函数，因此它可以传递给内核。 nvcc 在这里没有发出任何警告是正常的吗？
您已明确将其定义为主机和设备功能。内核启动在主机代码中运行。传递的结果函数引用是一个宿主函数

标签： c++ cuda functor

【解决方案1】：

Debian 不支持任何版本的 CUDA。 gcc 9.3 不是 CUDA 9.2 支持的工具

cuda 标签上有很多关于这些主题的问题。这个answer 链接到其中的一些。

简而言之，在主机代码中捕获__device__ 函数地址基本上是不可能的。内核启动（正如您在此处所拥有的）是用主机代码编写的；它是主机代码。因此say_hello的使用存在于主机代码中，它将引用__host__函数指针/地址。该函数指针/地址在设备代码中是无用的。（删除 __host__ 装饰器无济于事。）

有许多可能的解决方案，您已经探索过其中一种。传递包装在某种对象中的函数，当您直接使用 __device__ lambda 时，符合该描述。

另一种可能的解决您的函数指针方法不起作用的方法是在设备代码中捕获函数指针。然后它必须被传递到主机，然后它可以通过内核启动传递回设备代码，并在那里分派。上面的链接答案提供了多种实现方式。

【讨论】：