c++ - 如何将函数作为 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 标签上有很多问题涉及这些主题。这answer其中一些的链接。

简短的版本是根本不可能在主机代码中捕获 __device__ 函数地址。内核启动(如您在此处所用)是用主机代码编写的；它是主机代码。因此在主机代码中使用say_hello，它会引用__host__函数指针/地址。该函数指针/地址在设备代码中是无用的。 (删除 __host__ 装饰器将无济于事。)

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

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