如何使用 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 : 为虚拟计算架构编译时不允许选项‘--cubin (-cubin)’”

我正在使用 Tesla P100（计算 6.0），python3.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' )

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

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

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

How to use child kernels (CUDA dynamic parallelism) using PyCUDA

python-3.x

cuda

pycuda