使用 CUDA 模拟器调试 CUDA Python

CUDA 内置目标弃用通知

Numba 中内置的 CUDA 目标已被弃用，后续开发已移至 NVIDIA numba-cuda 包。请参阅 内置 CUDA 目标弃用和维护状态。

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。

模拟器的一些限制包括：

  • 它不执行类型检查/类型推断。如果 JIT 编译函数的任何参数类型不正确，或者任何局部变量的类型规范不正确，模拟器将不会检测到此问题。

  • 只模拟一个 GPU。

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

  • 大多数驱动程序 API 未实现。

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

  • Warp 和 Warp 级操作尚未实现。

  • 由于模拟器使用 Python 解释器执行内核，因此适用于硬件目标的按属性结构化数组访问可能会在模拟器中失败——请参阅 结构化数组访问。

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

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

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