调试#

打印值#

通常，最好的调试方法之一是简单地从内核中打印值。Warp 支持使用 print() 函数打印所有内置类型，例如：

v = wp.vec3(1.0, 2.0, 3.0)

print(v)

[1.0, 2.0, 3.0]

此外，对于标量类型，可以通过 wp.printf() 函数进行格式化的 C 风格打印，例如：

@wp.kernel
def mykernel():
    x = 1.0
    i = 2

    wp.printf("A float value %f, an int value: %d\n", x, i)

详细模式和打印启动#

对于复杂的应用程序，可能很难理解导致 bug 的操作顺序。为了帮助诊断这些问题，Warp 支持一个简单的选项，可以将所有启动和参数打印到控制台

wp.config.print_launches = True

详细模式也可以通过以下方式启用：

wp.config.verbose = True

在详细模式下，关于程序进度和代码生成（例如操作何时可能不可微）的附加消息将被打印到标准输出。

详细警告可以通过以下方式启用：

wp.config.verbose_warnings = True

这有助于识别特定的 Warp UserWarning 消息是从哪里发出的。

调试模式编译#

在调试模式下，Warp 内核将执行以下附加检查：

  • 如果存在超出定义形状的数组访问，则引发断言。

  • 如果 wp.tid() 在大网格上返回溢出值，则发出警告。

  • (仅限 GPU) 如果 CUDA 网格维度因块数溢出而被限制，则发出警告。

  • (仅限 GPU) 为设备代码生成行号信息。

在调试模式下启用 Warp 内核编译的最简单方法是设置：

wp.config.mode = "debug"

作为前述全局设置的替代方案，可以通过按模块设置以下内容来开启调试模式：

wp.set_module_options({"mode": "debug"})

断言#

assert 语句可以插入到 Warp 内核和用户定义函数中，以便在提供的布尔表达式评估为假时中断程序执行。断言仅在模块以调试模式编译时对其内核有效。

以下示例将在运行内核时引发断言，因为该模块是以调试模式编译的，并且 assert 语句期望传递给 expect_ones 内核的数组是一个全一数组，但我们传递了一个单元素全零数组：

import warp as wp

wp.config.mode = "debug"


@wp.kernel
def expect_ones(a: wp.array(dtype=int)):
    i = wp.tid()

    assert a[i] == 1, "Array element must be 1"


input_array = wp.zeros(1, dtype=int)

wp.launch(expect_ones, input_array.shape, inputs=[input_array])

wp.synchronize_device()

程序的输出将包含如下一行语句：

default_program:49: void expect_ones_133f9859_cuda_kernel_forward(wp::launch_bounds_t, wp::array_t<int>): block: [0,0,0], thread: [0,0,0] Assertion `("assert a[i] == 1, \"Array element must be 1\"",var_3)` failed.

单步调试#

可以将 Visual Studio 等 IDE 调试器附加到 Warp 进程，以便单步执行生成的内核代码。用户应首先通过设置以下内容在调试模式下编译内核：

wp.config.mode = "debug"

此设置确保正确生成行号和调试符号。启动 Python 进程后，应附加调试器并在生成的代码中插入断点。

注意

生成的内核代码与原始 Python 代码并非一一对应，但仍然可以重播单个操作并检查变量。

另请参阅 warp/tests/walkthrough_debug.py，了解如何在 CPU 上调试 Warp 内核代码的示例。

生成的代码#

有时，检查生成的代码以进行调试或性能分析可能会很有用。默认情况下，内核生成的代码存储在用户主目录的中心缓存位置。调用 wp.init() 时会在启动时打印缓存位置，例如：

Warp 0.8.1 initialized:
    CUDA Toolkit: 11.8, Driver: 11.8
    Devices:
    "cpu"    | AMD64 Family 25 Model 33 Stepping 0, AuthenticAMD
    "cuda:0" | NVIDIA GeForce RTX 3090 (sm_86)
    "cuda:1" | NVIDIA GeForce RTX 2080 Ti (sm_75)
    Kernel cache: C:\Users\LukasW\AppData\Local\NVIDIA Corporation\warp\Cache\0.8.1

内核缓存包含以 wp_ 开头的文件夹，其中包含生成的 C++/CUDA 代码以及运行时编译的每个模块的编译二进制文件。每个文件夹的名称都以由模块内容构建的十六进制哈希值结尾，以避免在使用多个进程时发生潜在冲突，并支持运行时定义内核的缓存。

如果怀疑 Warp 的内核缓存逻辑存在 bug，可以通过设置以下内容来禁用内核缓存：

wp.config.cache_kernels = True

CUDA 错误验证#

通过格式不正确的内核代码或输入，可能会产生越界内存访问违规。在这种情况下，CUDA 运行时将检测到违规并将 CUDA 上下文置于错误状态。随后的内核启动可能会静默失败，这可能导致难以诊断的问题。

如果怀疑存在 CUDA 错误，一个简单的验证方法是启用：

wp.config.verify_cuda = True

此设置将在每次 wp.launch() 后检查 CUDA 上下文，以确保其仍然有效。如果遇到错误，将引发异常，这通常有助于缩小问题内核的范围。

在捕获 CUDA 图时，不能使用 CUDA 错误验证。

注意

每次启动时验证 CUDA 状态需要同步 CPU 和 GPU，这会产生显着的开销。用户应确保仅在调试期间使用此设置。

检测非有限值#

wp.config.verify_fp = True 有助于识别计算何时产生非有限值，例如 NaN 或无穷大。单独使用此标志时，消息将被打印到标准输出流，指示正在检测无效值的函数。

如果与调试模式编译结合使用，则在检测到无效值时将引发断言。

CUDA 工具包调试工具#

Compute Sanitizer 工具（如 initcheck 和 memcheck）也可用于检测 Warp 应用程序中细微的内存访问问题，例如：

compute-sanitizer --tool initcheck python sim.py

Compute Sanitizer 套件可通过 CUDA 工具包获得。