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

Question

我想创建通用的 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 的消息。

Answer 1

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

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

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

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

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

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

How to pass a function as a cuda kernel parameter?

c++

cuda

functor