使用 CUDA 模拟器调试 CUDA Python

Numba 包含一个 CUDA 模拟器，它使用 Python 解释器和一些额外的 Python 代码实现了 CUDA Python 中的大部分语义。这可以用于调试 CUDA Python 代码，方法是在代码中添加打印语句，或者使用调试器逐步执行单个线程。

模拟器故意允许运行非CUDA代码，例如启动调试器和打印任意表达式以进行调试。因此，最好从为CUDA目标编译的代码开始，然后转移到模拟器以调查问题。

内核的执行由模拟器一次一个块地执行。为块中的每个线程生成一个线程，这些线程的调度由操作系统负责。

使用模拟器

通过在导入 Numba 之前将环境变量 NUMBA_ENABLE_CUDASIM 设置为 1，可以启用模拟器。然后可以正常执行 CUDA Python 代码。在核函数内部使用调试器的最简单方法是将单个线程停止，否则与调试器的交互将难以处理。例如，下面的核函数将在线程 <<<(3,0,0), (1, 0, 0)>>> 中停止:

@cuda.jit
def vec_add(A, B, out):
    x = cuda.threadIdx.x
    bx = cuda.blockIdx.x
    bdx = cuda.blockDim.x
    if x == 1 and bx == 3:
        from pdb import set_trace; set_trace()
    i = bx * bdx + x
    out[i] = A[i] + B[i]

当使用一维网格和一维块调用时。

支持的功能

该模拟器旨在尽可能完整地模拟在真实GPU上的执行——特别是，以下内容得到支持：

  • 原子操作

  • 常量内存

  • 本地内存

  • 共享内存：共享内存数组的声明必须放在单独的源代码行上，因为模拟器使用源代码行信息来跟踪跨线程的共享内存分配。

  • 映射数组。

  • 主机和设备内存操作：复制和设置内存。

  • syncthreads() 是被支持的 - 然而，在不同线程进入不同的 syncthreads() 调用的情况下，启动不会失败，但会出现意外行为。模拟器的未来版本可能会检测到这种情况。

  • 流API是支持的，但所有操作都是顺序且同步进行的，与在真实设备上不同。因此，在流上进行同步是一个无操作。

  • 事件API也得到支持，但不提供有意义的计时信息。

  • GPU 数据的传输 - 特别是使用 device_array() 和 device_array_like() 创建数组对象。固定内存的 API pinned() 和 pinned_array() 也被支持，但不会发生固定操作。

  • 支持 GPU 上下文列表（cuda.gpus 和 cuda.cudadrv.devices.gpus）的驱动 API 实现，并报告单个 GPU 上下文。此上下文可以像真实上下文一样关闭和重置。

  • 支持 detect() 函数，并报告一个名为 SIMULATOR 的设备。

  • 协作网格：可以启动一个协作内核，但只有一个块 - 模拟器总是从内核重载的 max_cooperative_grid_blocks() 方法返回 1。

模拟器的一些限制包括：

  • 它不执行类型检查/类型推断。如果对 jitted 函数的任何参数类型不正确，或者任何局部变量的类型指定不正确，模拟器将不会检测到这些错误。

  • 仅模拟了一个GPU。

  • 不支持对单个GPU的多线程访问，这会导致意外行为。

  • 大部分驱动API尚未实现。

  • 无法将 PTX 代码与 CUDA Python 函数链接。

  • 扭曲和扭曲级操作尚未实现。

  • 由于模拟器使用Python解释器执行内核，通过属性访问结构化数组在硬件目标上有效，但在模拟器中可能会失败 - 参见 结构化数组访问。

  • 直接对设备数组的操作仅部分支持，即，测试相等性、小于、大于以及基本数学运算是支持的，但许多其他操作，如原地操作符和位操作符则不支持。

  • 函数 ffs() 仅对可以使用32位整数表示的值正确工作。

显然，模拟器的速度也比实际设备的速度慢得多。为了使使用模拟器进行调试变得可行，可能需要减少输入数据的大小和CUDA网格的大小。