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

How to pass a function as a cuda kernel parameter?

本文关键字：CUDA 内核 参数传递 函数      更新时间：2023-10-16

我想创建一个通用的cuda内核，将可kerable对象作为参数(例如lambda或function(并调用。

我很难将设备函数传递给CUDA内核作为参数。

我的CUDA 9.2具有计算能力3.5。我在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的消息。

debian不是任何版本的CUDA的支持环境。GCC 9.3不是CUDA 9.2

的支持工具

在cuda标签上涵盖了这些主题有很多问题。此答案链接到其中许多。

简短的版本是从根本上不可能在主机代码中捕获__device__功能地址。内核启动(如您在这里所示(是用主机代码编写的；它是主机代码。因此，在主机代码中使用say_hello，它将参考__host__功能指针/地址。该功能指针/地址在设备代码中无用。(卸下__host__装饰器将无济于事。(

有许多可能的解决方案，您已经探索过其中之一。传递包裹在某种对象中的功能，并在直接按照您的身份直接使用__device__ lambda，适合该描述。

您所拥有的功能指针方法的另一个可能的修复程序是在设备代码中捕获功能指针。然后必须将其传递给主机，然后可以通过内核启动转换为设备代码，然后在此处派遣。上面的链接答案提供了多种可以实现的方法。

相关文章：