Skip to content

3.11。使用 CUDA 模拟器 调试 CUDA Python

原文： http://numba.pydata.org/numba-doc/latest/cuda/simulator.html

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

内核的执行由模拟器一次一个块执行。为块中的每个线程生成一个线程，并且将这些线程的执行调度留给操作系统。

3.11.1。使用模拟器

通过将环境变量 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]

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

3.11.2。支持的功能

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

  • 原子操作
  • 恒定记忆
  • 本地记忆
  • 共享内存：共享内存数组的声明必须位于不同的源代码行上，因为模拟器使用源代码行信息来跟踪跨线程的共享内存分配。
  • 支持 syncthreads() - 但是，在发散线程进入不同的 syncthreads() 调用的情况下，启动不会失败，但会发生意外行为。未来版本的模拟器可以检测到这种情况。
  • 支持流 API，但与实际设备不同，所有操作都按顺序和同步进行。因此，在流上进行同步是一种无操作。
  • 还支持事件 API，但不提供有意义的计时信息。
  • 与 GPU 之间的数据传输 - 特别是使用 device_array() 和 device_array_like() 创建数组对象。固定存储器 pinned() 和 pinned_array() 的 API 也受支持，但不会发生钉扎。
  • 支持 GPU 上下文列表（cuda.gpus和cuda.cudadrv.devices.gpus）的驱动程序 API 实现，并报告单个 GPU 上下文。这个上下文可以像真正的那样关闭和重置。
  • 支持 detect() 功能，并报告一个名为 SIMULATOR 的设备。

模拟器的一些限制包括：

  • 它不执行类型检查/类型推断。如果 jitted 函数的任何参数类型不正确，或者任何局部变量类型的规范不正确，模拟器将无法检测到。
  • 仅模拟一个 GPU。
  • 不支持对单个 GPU 的多线程访问，这将导致意外行为。
  • 大多数驱动程序 API 未实现。
  • 无法将 PTX 代码与 CUDA Python 函数链接。
  • Warp 和 warp 级操作尚未实现。

显然，模拟器的速度也远低于真实设备的速度。可能需要减小输入数据的大小和 CUDA 网格的大小，以便使模拟器易于调试。

回到顶部