如何使用 PyCUDA 使用子内核（CUDA 动态并行）

Question

我的 python 代码有一个 gpu 内核函数，它在主机的 for 循环中被多次调用，如下所示：

for i in range:   
    gpu_kernel_func(blocksize, grid)   

由于这个函数调用需要主机和gpu设备之间多次通信，效率不高，我想把它设为

gpu_kernel_function(){  
    for(){ 
        computation } ;  
}

但这需要额外的步骤来确保网格中的所有块都同步。根据动态并行性，调用虚拟子内核应确保每个线程（在整个网格中）都应在代码继续运行之前完成该子内核。所以我定义了另一个内核，就像 gpu_kernel_function 一样，我尝试了这个：

GPUcode = '''

\__global__ gpu_kernel_function() {... }  
\__global__ dummy_child_kernel(){ ... }
'''

gpu_kernel_function(){  
    for() {
        computation } ;  
    dummy_child_kernel(void);  
}

但我收到此错误“nvcc fatal : Option '--cubin (-cubin)' is not allowed whencompiled for a virtual compute architecture”

我正在使用 Tesla P100（计算 6.0）、python 3.5、cuda.8.0.44。我正在像这样编译我的源模块：

mod = SourceModule(GPUcode, options=['-rdc=true' ,'-lcudart','-lcudadevrt','--machine=64'],arch='compute_60' )

我也尝试了 compute_35，但出现了同样的错误。

Answer 1

错误消息明确告诉您问题所在。 compute_60 是 virtual architecture。您不能将虚拟架构静态编译为机器代码。它们旨在生成 PTX（虚拟机汇编器），以便由运行时将 JIT 转换为机器代码。 PyCUDA 使用 CUDA 工具链将代码编译为二进制有效负载（“cubin”），然后通过驱动程序 API 将其加载到 CUDA 上下文中。因此错误。

您可以通过指定有效的物理 GPU 目标架构来修复错误。所以你应该修改源模块构造函数调用如下：

mod = SourceModule(GPUcode, 
                   options=['-rdc=true','-lcudart','-lcudadevrt','--machine=64'],
                   arch='sm_60' )

这应该可以修复编译器错误。

但是，请注意，使用动态并行需要设备代码链接，我 99% 确定 PyCUDA 仍然不支持此功能，因此您可能无法通过 SourceModule 执行您所要求的操作.您可以使用 PyCUDA 外部的编译器手动链接您自己的 cubin，然后在 PyCUDA 中加载该 cubin。如果您搜索它们，您会发现许多如何正确编译动态并行性的示例。