CUDA-GDB

CUDA-GDB 用户手册，NVIDIA 工具，用于在 Linux 和 QNX 系统上调试 CUDA 应用程序。

1. 简介

本文档介绍了 CUDA-GDB，NVIDIA^® CUDA^® 调试器，用于 Linux 和 QNX 目标平台。

1.1. 什么是 CUDA-GDB？

CUDA-GDB 是 NVIDIA 工具，用于调试在 Linux 和 QNX 上运行的 CUDA 应用程序。CUDA-GDB 是 GDB（GNU 项目调试器）的扩展。该工具为开发人员提供了一种在实际硬件上调试 CUDA 应用程序的机制。这使开发人员能够调试应用程序，而无需考虑仿真和模拟环境可能引入的潜在差异。

1.2. 支持的功能

CUDA-GDB 旨在为用户提供无缝的调试环境，允许在同一应用程序中同时调试 GPU 和 CPU 代码。正如 CUDA C 编程是 C 编程的扩展一样，使用 CUDA-GDB 进行调试是使用 GDB 进行调试的自然扩展。现有的 GDB 调试功能本身就存在于调试主机代码中，并且还提供了其他功能来支持调试 CUDA 设备代码。

CUDA-GDB 支持调试 C/C++ 和 Fortran CUDA 应用程序。Fortran 调试支持仅限于 64 位 Linux 操作系统。

CUDA-GDB 允许用户设置断点、单步执行 CUDA 应用程序，还可以检查和修改硬件上运行的任何给定线程的内存和变量。

CUDA-GDB 支持调试所有 CUDA 应用程序，无论它们是使用 CUDA 驱动程序 API、CUDA 运行时 API 还是两者都使用。

CUDA-GDB 支持调试为特定 CUDA 架构（例如 sm_75 或 sm_80）编译的内核，但也支持调试在运行时编译的内核，称为即时编译（just-in-time compilation），简称 JIT 编译。

1.3. 关于本文档

本文档是 CUDA-GDB 的主要文档，其组织形式更像是用户手册而不是参考手册。文档的其余部分将介绍如何安装和使用 CUDA-GDB 调试 CUDA 内核，以及如何使用添加到 GDB 的新 CUDA 命令。还提供了一些演练示例。假定用户已经了解用于调试主机应用程序的基本 GDB 命令。

2. 发行说明

12.8 发行版

Maxwell、Pascal 和 Volta 弃用通知

• 对 Maxwell、Pascal 和 Volta 设备（sm 5.x、sm 6.x 和 sm 7.x）的支持已弃用。支持将在未来的版本中删除。

核心转储环境变量弃用通知

• CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION 和 CUDA_ENABLE_LIGHTWEIGHT_COREDUMP 环境变量在 CUDA 12.5 中已被弃用，取而代之的是 CUDA_COREDUMP_GENERATION_FLAGS。

功能

• 添加了解析 JSON cuobjdump 输出以进行 CUDA 反汇编的功能。

• 添加了使用新的 maint reset_cuda_stats 命令重置 CUDA 统计信息收集指标的功能。

• 将 set/show cuda collect_stats 重命名为 maint set/show cuda_stats。

• 添加了对处理 CUDA 堆栈金丝雀异常的支持。

• 通过删除对跟踪上下文推送/弹出事件的支持来提高性能。

• 增加了对 QNX SDP 8.0 的支持。

• 增加了对 __nv_fp8_e8m0 类型的支持。

修复的问题

• 修复 gdb/mi 输出记录，使其符合 CUDA 焦点更改消息的输出记录格式。

• 修复了加载没有有效上下文的 CUDA 核心转储文件的问题。

• 修复了在调试 CUDA 核心转储文件时，从 CUDA 焦点切换到主机焦点时崩溃的问题。

• 修复了在调试 CUDA 核心转储文件时使用 info threads 时崩溃的问题。

12.7 发行版

功能

• 现在可以通过命令 generate-core-file 按需生成非致命核心转储文件。

• 在生成 CUDA 反汇编时，使用 posix_spawn 调用 cuobjdump。

• 为各种 CUDA 概念添加了 Python 扩展。

• CUDA 延迟加载的性能改进。

• 在可用时，打印驱动程序/cudart 错误的扩展错误字符串。

• 添加了在遇到集群异常时识别目标块的支持。

• 添加了通过新的内置 clusterDim 变量查询首选集群大小的功能。

修复的问题

• 修复了在遇到无效 CUDA 寄存器号时，由于 CUDA 寄存器外推失败而导致的崩溃。

• 修复了阻止切换到具有活动 CUDA inferior 的 cudacore 目标的问题。

• 添加了对 OptiX 编译器生成的 DW_OP_bregx 中 ASCII regno 的识别。

12.6 发行版

功能

• 新的 set/show cuda step_divergent_lanes 命令，用于控制在聚焦的 CUDA 线程不再活动时自动单步执行发散线程。默认情况下处于启用状态，以匹配现有行为。

• 提高了打开包含许多 cubin 的 GPU 核心文件的性能。

• 允许用户使用 ctrl-c 并中断 GPU 核心文件的加载。

• 当开始单步执行发散线程且 step_divergent_lanes 处于启用状态时，发出警告。

• 当从 CUDA 线程切换回主机线程焦点时，发出线程切换通知。

修复的问题

• 修复了当当前线程焦点超出范围时，切换到最近邻线程时 CUDA 线程选择的问题。

• 修复了调试 GPU 核心文件时打印 CUDA 寄存器的问题。

• 增强了在遇到具有损坏的扩展行信息的 cubin 时的错误处理。

• 修复了接收 QNX cuda-gdbserver 的乱序数据包时发生的错误。

12.5 发行版

更新的 GDB 版本

• 从 GDB 13.1 升级到 13.2。请参阅 GDB 13.2 更改

支持移除通知

• 已移除对 CUDA-GDB 的 macOS 主机客户端的支持。

• 已移除对 Android 的支持。

• 已移除对 Python 3.6 和 3.7 的支持。

功能

• 多构建功能，支持所有受支持平台上的原生 Python 和 TUI 模式。cuda-gdb 程序现在是一个包装脚本，用于调用适当的 cuda-gdb 二进制文件。如果未检测到受支持的 Python 或 libncurses，则包装器将回退到禁用 Python 和 TUI 支持的 cuda-gdb 二进制文件。

• 增加了对 TUI 模式的支持。

• 增加了对 Python 3.10、3.11 和 3.12 的支持。

• 增加了对检测和打印退出 warp 中遇到的异常的支持。这可能发生在调试启用了优化的应用程序时。

• 为 info cuda managed 和 info cuda line 添加了新的 gdb/mi 命令等效项。

修复的问题

• 修复了打印 CUDA 函数的引用参数实参的问题。

• 修复了从/向 CUDA 通用内存读取/写入时导致崩溃/错误的问题。

• 修复了对于同一内核的背靠背启动，break_on_launch 断点被遗漏的问题。

• 修复了当断点在发散线程中命中时，错误地将断点命中事件报告为 SIGTRAP 的问题。

• 修复了 QNX 上 cuda-gdbserver 数据包乱序到达时崩溃的问题。

• 改进了在读取 CUDA 反汇编时遇到错误时的错误处理。

• 改进了从致命 CUDA 异常恢复执行时的退出处理。

12.4 发行版

更新的 GDB 版本

• 从 GDB 12.1 升级到 13.1。请参阅 GDB 13.1 更改

Android 弃用通知

• 已弃用对 Android 的支持。它将在未来的版本中删除。

Python 3.6 和 3.7 弃用通知

• 已弃用对生命周期结束的 Python 3.6 和 3.7 版本的支持。它将在未来的版本中删除。

功能

• 性能增强，减少了 CUDA 调试器 API 调用的总数。

• 加载具有大量 GPU 寄存器的大型 cubin 的性能增强。

• 单步执行 warp 范围屏障时的性能增强。

• 增加了从 GPU 核心转储打印常量存储体中包含的值的支持。

修复的问题

• 防止了在反汇编时对克隆函数名称进行 shell 扩展。

• 修复了在未知符号名称上设置条件断点时崩溃的问题。

• 修复了在全球指针上设置观察点的问题。

• 修复了 inferior 拆卸期间 switch_to_thread_1 中的断言。

• 修复了较新的 Intel 处理器遇到的附加失败问题。

• 重构了 libpython 层，以避免不必要的 gdb 代码更改。

12.3 发行版

macOS 主机客户端弃用通知

• 已弃用对 CUDA-GDB 的 macOS 主机客户端的支持。它将在未来的版本中删除。

功能

• 增加了打印常量存储体中包含的值的支持。新的 $_cuda_const_bank(bank, offset) 便利函数，用于获取常量存储体中偏移量的地址。请参阅 常量存储体。

• 添加了性能增强功能，可减少运行具有许多 CUDA 线程的应用程序时的开销。

• 增加了对 CUDA 函数指针的支持。

修复的问题

• 修复了从附加进程分离时可能导致崩溃的问题。

• 修复了几个 info cuda 命令中存在的线程排序问题。

• 增加了在设备上不存在有效 warp 时打开 GPU 核心转储文件的支持。

• 添加了 OptiX 使用的缺失 DWARF 运算符。

• 修复了解析 CUDA Fortran 指针类型的问题。

• 修复了在不存在 CUDA 集群时显示 CUDA 集群坐标的问题。

12.2 发行版

功能

• 启用了在遇到 CUDA 调试器 API 错误时打印扩展错误消息的功能。

• 启用了在具有开发工具模式的机密计算模式下进行调试的支持。请参阅 机密计算部署指南 <https://docs.nvda.net.cn/confidential-computing-deployment-guide.pdf>，了解有关如何启用该模式的更多详细信息。

修复的问题

• 修复了 OptiX 应用程序中回溯中出现的“??”。

• 现在在启用 CUDA 延迟加载的情况下，可以正确处理主机影子断点。

• 修复了调试 LLVM 生成的 cubin 时出现的名称修饰问题。

• 现在可以正确显示 CUDA 集群坐标。

• 修复了在使用 cuda-gdbserver 进行远程调试时，附加到使用 CUDA 延迟加载的应用程序的问题。

12.1 发行版

添加了用于控制核心转储行为的 CUDA 驱动程序 API

• CTK 12.1 和 r530 驱动程序添加了新的 API，允许开发人员在其应用程序内部以编程方式启用/配置核心转储设置，而不是使用环境变量。有关更多信息，请参阅 CUDA 驱动程序 API 手册。

功能

• 使用 CUDA 延迟加载的应用程序的性能改进。

• 增加了对具有大量节（超过 32767 个）的 ELF cubin 的支持。

• 为 CUDA Graphs 增加了 break_on_launch 支持。

修复的问题

• 移除了不受支持的 set/show gpu_busy_check 命令。

• 在 QNX 上，修复了 info threads 错误报告死主机线程的问题。

• 修复了单步执行/下一步跳过内联函数调用的性能问题。

• 修复了使用 info cuda managed 命令时的性能问题。

• 修复了使用 set follow-fork-mode child 时的问题。

• 修复了解析 DWARF 以进行自引用结构的问题。

12.0 发行版

更新的 GDB 版本

• 从 GDB 10.2 升级到 12.1。请参阅 GDB 12.1 更改

纹理和表面引用支持已移除

• CTK 12.0 移除了对纹理和表面引用 API 的支持。已移除对打印纹理和表面引用的支持。

CUDA 内存检查器集成已移除

• cuda-memcheck 已在 CUDA 11.x 中弃用，并由 Compute Sanitizer 取代。新的内存检查工作流程是从 CLI 使用 Compute Sanitizer。这将支持在检测到问题时生成核心转储文件，然后可以使用 CUDA-GDB 打开和检查这些文件，类似于其他核心转储文件。CUDA 12.0 发行版已移除对 cuda-memcheck 的支持。

调试使用 CUDA 动态并行性的应用程序

• 默认情况下，对于使用 CTK 12.0 或更高版本编译的应用程序，已移除对使用经典调试器后端或在 Maxwell GPU 上调试使用 CUDA 动态并行性的应用程序的支持。在这些情况下，可以通过重新编译应用程序并传递 -DCUDA_FORCE_CDP1_IF_SUPPORTED 标志来完成调试。

功能

• 从基本 gdb/10.2 升级到 gdb/12.1。

• 增加了对线程块集群的初始支持。

• 将 --cuda-use-lockfile 的默认行为更改为 0。默认情况下不再创建锁定文件。

修复的问题

• 解决了在单步执行设备系统调用时可能遇到的挂起问题。

• 修复了在 info cuda 命令中显示活动 warp 掩码时的溢出问题。

• 更改了内部 CUDA 动态并行性检测断点，使其仅在启用 break_on_launch 时设置。

• 移除了不受支持的 gpu_busy_check setting。

11.8 发行版

功能

• 默认情况下使用新的统一调试器 (UD) 调试后端。

• 增加了对调试使用 CUDA 延迟加载的应用程序的支持。

• 现在在 Windows Subsystem for Linux (WSL) 上启用了调试器。

• 增加了打印 FP8 值（E4M3 和 E5M2）的基本类型支持。

注释

• 默认情况下，CUDA-GDB 将使用新的统一调试器 (UD) 后端。此更改对于大多数使用 Pascal 或更新版本显卡的用户是透明的。对于 Maxwell 调试，或强制使用旧的经典调试后端，请在您的环境中将 CUDBG_USE_LEGACY_DEBUGGER 设置为 1。

• 在此版本中，GH100 平台不支持 WSL。

11.7 发行版

功能

• 主要的 break_on_launch 性能增强，使用新的 KERNEL_READY 通知机制而不是设置手动断点。

• 重构了 info cuda 命令输出，使其更加简洁。省略了打印非活动消息。

• 添加了新的 --disable-python 命令行选项，以禁用 Python 解释器 dlopen。

修复的问题

• 修复了 follow-fork child，以避免父进程和子进程都使用 CUDA 时出现挂起行为。

• 添加了 libpython 函数的缺失 dlsym，该函数导致某些版本的 libpython 出现错误。

11.6 发行版

更新的 GDB 版本

• 从 GDB 10.1 升级到 10.2。请参阅 GDB 10.2 更改

功能

• 在反汇编视图中添加了 errorpc 指令前缀。如果设置了错误 PC，则以 *> 为指令添加前缀。

修复的问题

• 修复了 lineinfo 帧以正确显示源文件名。

• 修复了写入从主机分配的 gpu 全局内存的问题。

• 修复了阻止在某些情况下读取主机变量的错误。

• 修复了阻止 QNX 启动的 cuda-gdbserver 初始化检查。

11.5 发行版

Jetson 和 Drive Tegra 设备上的 Python 3 支持

• 已移除对 Python 2 的支持。CUDA-GDB 现在支持 Jetson 和 Drive Tegra 设备上的 Python 3。

修复的问题

• 在动态加载 libpython3 库时，添加了强大的版本检查。加载的 libpython3 将与 PATH 中的 python3 运行时的版本匹配。

• 添加了在加载 libpython3 库时检查 PEP-3149 标志名称的支持。

• 增加了对 Python 3.9 的动态加载的支持。

• 修复了在某些 RHEL 发行版上覆盖 PYTHONPATH 的问题。

11.4 Update 1 发行版

Fedora 34 的已知问题

• CUDA-GDB 在 Fedora 34 上进行调试时存在已知问题，可能不可靠。

修复的问题

• 为 ppc64le 和 aarch64 SBSA 启用了 python 集成。

• 修复了调试 CUDA 应用程序时的性能回归。

• 修复了通过 cuda-gdbserver 进行远程调试时出现的间歇性挂起问题。

• 修复了 set cuda api_failures stop 不会触发失败断点的错误。

• 更改了 python 行为，以 dlopen 与 PATH 中的 python3 解释器版本匹配的 libpython 库。

• OpenMP Fortran：修复了在 OpenMP 并行区域内设置断点时崩溃的问题。

• OpenMP：更好地支持打印并行区域内的局部变量。

• Fortran：增加了对打印假定形状数组和数组切片的更新支持。

• 修复了 cudacore 调试中在主机和设备线程焦点之间进行选择的问题。

• QNX 远程调试的各种修复。

11.4 发行版

更新的 GDB 版本

• 从 GDB 8.3 升级到 10（基于 GDB 10.1）。请参阅 GDB 10.1 更改

Python 3 支持

• 已移除对 Python 2 的支持。CUDA-GDB 现在支持 Python 3。

GDB TUI 模式已禁用

• 已禁用对 GDB TUI 模式的支持。这避免了缺少 ncurses-5.5 支持的操作系统出现跨平台依赖项不匹配的情况。

Kepler 弃用通知

• 已弃用对 Kepler 设备（sm_35 和 sm_37）的支持。Kepler 支持将在未来的版本中删除。

核心转储支持

• 增加了使用 CUDA_COREDUMP_FILE 将核心转储文件写入命名管道的支持。

修复的问题

• 增加了在核心转储文件中显示 SIGTRAP 异常的支持。

• 禁用了在调试 CUDA 目标时启用调度程序锁定的功能。

• 修复了 cuda_register_name 和 cuda_special_register_name 以避免在错误时返回旧的缓存结果。

• 修复了创建 CUDA 临时目录时的间歇性竞争条件。

• QNX 远程调试的各种修复。

11.3 发行版

Python 2 弃用通知

• 即将弃用对 Python 2 的支持。CUDA-GDB 将在未来的版本中迁移到使用 Python 3 支持进行构建。

修复的问题

• 改进了远程调试的延迟附加。

11.2 Update 1 发行版

GDB TUI 弃用通知

• 即将弃用对 GDB TUI 模式的支持。这将避免缺少 ncurses-5.5 支持的操作系统出现跨平台依赖项不匹配的情况。GDB TUI 模式将在未来的版本中禁用。

修复的问题

• 修复了在运行 CPU 代码时打印全局 GPU 内存中的字符串的问题。

• 修复了扩展 debug_line 处理中的错误。

• 修复了内置 gdb 变量（如 gridDim）的截断问题。

• 修复了 DWARF dies 缺少名称时启动期间发生的段错误。

• 修复了 CUDA 内核调用 assert 时发生的段错误。

• 修复了阻止调试 > 2GB cubin 的错误。

• 为使用 --lineinfo 编译的 cubin 添加了小的可用性增强功能。

• 修复了在 CLion 中使用 CUDA-GDB 时由漂亮的打印机引起的段错误。

11.1 发行版

更新的 GDB 版本

• 从 GDB 8.2 升级到 8.3（基于 gdb 8.3.1）。请参阅 gdb 8.3.1 更改

对 SM 8.6 的支持

• CUDA-GDB 现在支持计算能力为 8.6 的设备。

更新的 DWARF 解析器

• 旧的二进制文件可能需要重新编译，以确保 CUDA 特定的 DWARF 信息是最新的。

修复的问题

• 修复了附加到正在运行的 CUDA 进程时的间歇性死锁。

• 修复了检查半精度寄存器值时的错误。

11.0 发行版

更新的 GDB 版本

• CUDA-GDB 已从 GDB/7.12 升级到 GDB/8.2。

对 SM8.0 的支持

• CUDA-GDB 现在支持计算能力为 8.0 的设备。

对 Bfloat16 的支持

• 增加了对 Bfloat16 (__nv_bfloat16) 类型的支持。

MIG 支持

• CUDA-GDB 支持 MIG。每个 MIG 实例上可以有一个单独的调试器会话。如果需要多个调试器，请参阅 多个调试器。

Mac 支持

• 不再支持在 macOS 上进行调试。但是，macOS 仍然可以用作主机系统（其中 CUDA-GDB 在 macOS 下运行，使用 cuda-gdbserver 调试远程目标）。macOS 版本的 CUDA-GDB 下载可以在以下位置找到：在此处下载

10.1 发行版

增强了仅使用行号信息进行调试的功能

• 对 CUDA-GDB 支持调试使用 -lineinfo 但未使用 -G 编译的程序进行了多项增强。这主要用于调试使用 OptiX/RTCore 构建的程序。另请参阅 使用行号信息编译。

10.0 发行版

Turing 统一寄存器支持

• 增加了对检查和修改 Turing GPU 上的统一寄存器的支持。

9.2 发行版

用户引发的核心转储支持

• 对于支持计算抢占的设备，增加了用户引发的核心转储支持。可以使用新的环境变量：CUDA_ENABLE_USER_TRIGGERED_COREDUMP 来启用此功能。

9.1 发行版

Volta-MPS 核心转储支持

• Volta-MPS 上支持 GPU 核心转储生成。

轻量级 GPU 核心转储支持

• CUDA-GDB 支持读取轻量级 GPU 核心转储文件。可以使用新的环境变量：CUDA_ENABLE_LIGHTWEIGHT_COREDUMP 来启用此功能。

7.0 发行版

GPU 核心转储支持

• CUDA-GDB 支持读取 GPU 和 GPU+CPU 核心转储文件。可以使用新的环境变量：CUDA_ENABLE_COREDUMP_ON_EXCEPTION、CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION 和 CUDA_COREDUMP_FILE 来启用和配置此功能。

6.5 发行版

CUDA Fortran 支持

• CUDA-GDB 支持在 64 位 Linux 操作系统上进行 CUDA Fortran 调试。

GDB 7.6.2 代码库

• CUDA-GDB 的代码库已升级到 GDB 7.6.2。

6.0 发行版

统一内存支持

• 托管变量可以从主机线程或设备线程读取和写入。调试器还会使用 @managed 注释位于托管内存中的内存地址。可以通过新的 info cuda managed 命令访问静态分配的托管变量列表。

GDB 7.6 代码库

• CUDA-GDB 的代码库已从 GDB 7.2 升级到 GDB 7.6。

Android 支持

• CUDA-GDB 现在可以用于本地或远程调试 Android 原生应用程序。

单步执行优化

• CUDA-GDB 现在可以使用优化的方法来单步执行程序，这在大多数情况下可以加速单步执行。可以通过执行 set cuda single_stepping_optimizations off 来禁用此功能。

更快的远程调试

• 为了使远程调试速度显著提高，高达 2 个数量级，我们付出了很多努力。这些努力也使本地调试速度更快。

内核入口断点

• set cuda break_on_launch 选项现在将在从 GPU 启动的内核上中断。此外，启用此选项不会影响内核启动通知。

精确的错误归因

• 在 Maxwell 架构 (SM 5.0) 上，将准确报告触发异常的指令。应用程序会持续向前执行，调试器停止的 PC 可能与该地址不匹配，但额外的输出消息会标识异常的来源。

生命周期优化

• 为了缓解在某些代码地址处变量不可访问的问题，调试器提供了两个新选项。使用 set cuda value_extrapolation，最新已知值将以 (possibly) 前缀显示。使用 set cuda ptx_cache，与源变量关联的 PTX 寄存器的最新已知值将以 (cached) 前缀显示。

事件通知

• 内核事件通知默认情况下不再显示。

• 添加了新的内核事件详细程度选项：set cuda kernel_events, set cuda kernel_events_depth。此外，set cuda defer_kernel_launch_notifications 已被弃用，不再起作用。

5.5 版本

内核启动跟踪

• 引入了两个新命令 info cuda launch trace 和 info cuda launch children，用于在使用动态并行时显示内核启动跟踪和给定内核的子内核。

单 GPU 调试 (BETA)

• CUDA-GDB 现在可以用于调试在渲染桌面 GUI 的同一 GPU 上的 CUDA 应用程序。此功能还支持调试长时间运行或无限期的 CUDA 内核，否则这些内核会遇到启动超时。此外，多个 CUDA-GDB 会话可以在同一 GPU 上调试上下文切换的 CUDA 应用程序。此功能在 Linux 上适用于 SM3.5 设备。有关启用此功能的信息，请参阅 Single-GPU Debugging with the Desktop Manager Running 和 Multiple Debuggers。

远程 GPU 调试

• CUDA-GDB 与 CUDA-GDBSERVER 结合使用，现在可以用于调试在远程主机上运行的 CUDA 应用程序。

5.0 版本

动态并行支持

• CUDA-GDB 完全支持动态并行，这是 5.0 工具包中引入的一项新功能。调试器能够跟踪从另一个内核启动的内核，并像检查和修改任何其他 CPU 启动的内核一样检查和修改变量。

附加/分离

• 现在可以附加到已在运行的 CUDA 应用程序。也可以在让应用程序运行完成之前从应用程序分离。附加后，调试器的所有常用功能都可供用户使用，就像应用程序是从调试器启动的一样。使用动态并行的应用程序也支持此功能。

异常时附加

• 使用环境变量 CUDA_DEVICE_WAITS_ON_EXCEPTION，应用程序将正常运行，直到发生设备异常。然后应用程序将等待调试器附加到它以进行进一步调试。

API 错误报告

• 检查所有 CUDA 驱动程序 API 和 CUDA 运行时 API 函数调用的错误代码对于确保 CUDA 应用程序的正确性至关重要。现在，当任何 API 调用返回错误时，调试器能够报告甚至停止。有关更多信息，请参阅 set cuda api_failures。

内联子例程支持

• 现在可以从 SM 2.0 及更高版本的调试器访问内联子例程。用户可以检查这些子例程的局部变量，并访问调用帧堆栈，就像这些例程不是内联的一样。

4.2 版本

Kepler 支持

• CUDA-GDB 4.2 版本的主要更改是增加了对新 Kepler 架构的支持。此版本中没有其他用户可见的更改。

4.1 版本

源代码库升级到 GDB 7.2

• 到目前为止，CUDA-GDB 在 Linux 上是基于 GDB 6.6，在 Darwin (Apple 分支) 上是基于 GDB 6.3.5。现在，两个版本的 CUDA-GDB 都使用相同的 7.2 源代码库。

• 现在 CUDA-GDB 支持更新版本的 GCC（已测试到 GCC 4.5），更好地支持 DWARF3 调试信息，并具有更好的 C++ 调试支持。

并发会话支持

• 在 4.1 版本中，取消了单个 CUDA-GDB 进程的限制。现在，只要正在处理的应用程序之间不共享 GPU，就允许共存多个 CUDA-GDB 会话。例如，一个 CUDA-GDB 进程可以使用 GPU 0 调试进程 foo，而另一个 CUDA-GDB 进程可以使用 GPU 1 调试进程 bar。可以使用 CUDA_VISIBLE_DEVICES 环境变量强制独占使用 GPU。

新的 'autostep' 命令

• 添加了一个新的 'autostep' 命令。该命令通过自动单步执行代码段来提高 CUDA 异常的精度。

• 在正常执行情况下，可能无法精确报告发生异常的线程和指令。但是，如果在发生异常时程序正在单步执行，则可以确定生成异常的确切指令。

• 手动单步执行程序是一个缓慢而乏味的过程。因此，'autostep' 通过允许用户指定他们怀疑可能发生异常的代码段来帮助用户。当程序运行时，这些代码段会自动单步执行，并且会精确报告在这些代码段中发生的任何异常。

• 从 CUDA-GDB 输入 'help autostep' 以获取该命令的语法和用法。

多上下文支持

• 在计算能力为 SM20 或更高的 GPU 上，现在支持在同一 GPU 上调试多个上下文。这在以前是一个已知的限制。

设备断言支持

• 随 4.1 版本工具包发布的 R285 驱动程序支持设备断言。CUDA-GDB 支持断言调用，并在命中断言时停止应用程序的执行。然后可以像往常一样检查变量和内存。如果需要，应用程序也可以在断言之后恢复。使用 'set cuda hide_internal_frames' 选项来显示/隐藏系统调用帧（默认情况下隐藏）。

临时目录

• 默认情况下，调试器 API 将使用 /tmp 作为存储临时文件的目录。要选择其他目录，必须设置 $TMPDIR 环境变量和 API CUDBG_APICLIENT_PID 变量。

3. 入门

可以按照 快速入门指南 中的说明安装 CUDA 工具包。

应采取进一步的步骤来设置调试器环境、构建应用程序和运行调试器。

3.1. 设置调试器环境

3.1.1. 临时目录

默认情况下，CUDA-GDB 使用 /tmp 作为存储临时文件的目录。要选择其他目录，请设置 $TMPDIR 环境变量。

注意

用户必须对 CUDA-GDB 使用的临时目录具有写入和执行权限。否则，调试器将因内部错误而失败。

注意

$TMPDIR 的值在应用程序和 CUDA-GDB 的环境中必须相同。如果它们不匹配，CUDA-GDB 将无法附加到应用程序进程。

注意

由于 Android 设备上不存在 /tmp 文件夹，因此必须在启动 cuda-gdb 之前设置 $TMPDIR 环境变量并将其指向用户可写入的文件夹。

3.1.2. 在 Jetson 和 Drive Tegra 设备上使用 CUDA-GDB 调试器

默认情况下，在 Jetson 和 Drive Tegra 设备上，仅当 cuda-gdb 和 cuda-gdbserver 由 debug 组的成员用户启动时，才支持 GPU 调试。

要将当前用户添加到 debug 组，请运行以下命令

sudo usermod -a -G debug $USER

3.2. 编译应用程序

3.2.1. 调试编译

NVCC，NVIDIA CUDA 编译器驱动程序，提供了一种机制来生成 CUDA-GDB 正常工作所需的调试信息。当编译应用程序以便于使用 CUDA-GDB 进行调试时，必须将 -g -G 选项对传递给 NVCC；例如，

nvcc -g -G foo.cu -o foo

使用此行编译 CUDA 应用程序 foo.cu

• 强制进行 -O0 编译，但非常有限的死代码消除和寄存器溢出优化除外。

• 使编译器在可执行文件中包含调试信息

注意

启用 -G 选项会通过包含调试信息来增加二进制文件大小，并由于缺少编译器优化而降低性能。

要使用 CUDA-GDB 正常工作所需的调试信息编译 CUDA Fortran 代码，必须使用 -g 选项调用 pgfortran，PGI CUDA Fortran 编译器。此外，为了便于调试以及与未来 GPU 架构的向前兼容性，建议使用 -Mcuda=nordc 选项编译代码；例如，

pgfortran -g -Mcuda=nordc foo.cuf -o foo

有关可用编译标志的更多信息，请查阅 PGI 编译器文档。

3.2.2. 使用行号信息进行编译

对 cuda-gdb 的支持进行了多项增强，以调试使用 -lineinfo 而不是 -G 编译的程序。这主要用于调试使用 OptiX/RTCore 构建的程序。

请注意，尝试调试优化代码时可以使用 -lineinfo。在这种情况下，调试器单步执行和断点行为可能显得有些不稳定。

• 单步执行时，PC 可能会意外地向前和向后跳转。

• 用户可能会单步执行到没有行号信息的代码中，从而导致无法确定 PC 处的代码属于哪个源文件/行号。

• 断点可能会在与最初设置的行不同的行上中断。

调试 OptiX/RTCore 代码时，应牢记以下几点

• 用户无法调试或检查 NVIDIA 内部代码。

• OptiX/RTCode 调试仅限于 -lineinfo，不支持使用完整调试信息 (-G) 构建此代码。

• OptiX/RTCode 代码是高度优化的，因此上述关于调试优化代码的注意事项也适用。

3.2.3. 使用 PTX 调试选项进行编译

使用 -Xptxas nvcc 选项，可以使用 --make-errors-visible-at-exit 进行错误检测。

--make-errors-visible-at-exit                                          (-make-errors-visible-at-exit)
     Generate required instructions at exit point to make memory faults and errors visible at exit.

注意

此标志生成的额外指令可能会降低应用程序性能。

3.2.4. 为特定 GPU 架构编译

默认情况下，编译器将仅为 compute_52 PTX 和 sm_52 cubin 生成代码。对于较新的 GPU，内核将在运行时从 PTX 重新编译，以适应目标 GPU 的架构。为特定的虚拟架构编译可保证应用程序适用于之后的任何 GPU 架构，但会牺牲性能。这样做是为了向前兼容。

强烈建议为应用程序的目标 GPU 架构一次性编译应用程序，并为最新的虚拟架构生成 PTX 代码，以实现向前兼容。

GPU 架构由其计算能力定义。有关 GPU 列表及其各自的计算能力，请参阅 https://developer.nvidia.com/cuda-gpus。同一应用程序可以为多个 GPU 架构编译。使用 -gencode 编译选项来指定要为哪个 GPU 架构编译。该选项可以多次指定。

例如，要为计算能力为 7.0 的 GPU 编译应用程序，请将以下标志添加到编译命令

-gencode arch=compute_70,code=sm_70

要为计算能力超过 7.0 的任何未来架构编译 PTX 代码，请将以下标志添加到编译命令

-gencode arch=compute_70,code=compute_70

有关更多信息，请查阅编译器文档，网址为 https://docs.nvda.net.cn/cuda/cuda-compiler-driver-nvcc/index.html#extended-notation

3.3. 使用调试器

CUDA-GDB 可以在以下系统配置中使用

3.3.1. 在桌面管理器运行时进行单 GPU 调试

对于计算能力为 6.0 及更高的设备，CUDA-GDB 可用于调试在运行桌面 GUI 的同一 GPU 上的 CUDA 应用程序。

此外，对于计算能力低于 6.0 的设备，可以使用软件抢占来调试在运行桌面 GUI 的同一 GPU 上的 CUDA 应用程序。有两种方法可以启用此功能

注意

这是一个在 Linux 上可用的 BETA 功能，仅在 Maxwell 上受支持。对于计算能力为 SM6.0 及更高的 GPU，以下列出的选项将被忽略。

• 使用以下命令

  set cuda software_preemption on
  

• 导出以下环境变量

  CUDA_DEBUGGER_SOFTWARE_PREEMPTION=1
  

以上任一选项都将激活软件抢占。必须在运行应用程序之前设置这些选项。当 GPU 命中断点或任何其他通常会导致 GPU 冻结的事件时，CUDA-GDB 会释放 GPU 以供桌面或其他应用程序使用。这使 CUDA-GDB 能够在运行桌面 GUI 的同一 GPU 上调试 CUDA 应用程序，并支持在同一 GPU 上调试多个 CUDA 应用程序上下文切换。

3.3.2. 多 GPU 调试

多 GPU 调试是指应用程序在多个支持 CUDA 的设备上运行的情况。多 GPU 调试与单 GPU 调试没有太大区别，除了几个额外的 CUDA-GDB 命令，可让您在 GPU 之间切换。

任何 GPU 命中断点都将暂停在该系统上运行 CUDA 的所有 GPU。暂停后，您可以使用 info cuda kernels 查看所有活动内核及其运行所在的 GPU。当任何 GPU 恢复时，所有 GPU 都会恢复。

注意

如果使用 CUDA_VISIBLE_DEVICES 环境，则只会暂停和恢复指定的设备。

所有支持 CUDA 的 GPU 都可以运行一个或多个内核。要切换到活动内核，请使用 cuda kernel <n>，其中 n 是从 info cuda kernels 检索的内核 ID。

注意

同一个内核可以由不同的上下文和设备同时加载和使用。当在此类内核中设置断点时（通过名称或文件名和行号），它将任意解析为该内核的唯一实例。使用运行时 API，无法控制断点将解析为的确切实例。使用驱动程序 API，用户可以通过在模块加载后立即设置断点来控制断点将解析为的实例。

3.3.3. 远程调试

有多种方法可以使用 CUDA-GDB 远程调试应用程序。除了使用主机系统上的 SSH 或 VNC 连接到目标系统外，还可以使用 target remote GDB 功能。使用此选项，本地 cuda-gdb（客户端）连接到在目标系统上运行的 cuda-gdbserver 进程（服务器）。Linux 客户端和 Linux 或 QNX 服务器支持此选项。

以这种方式设置远程调试是一个 2 步过程

在远程主机上启动 cuda-gdbserver

cuda-gdbserver 可以在远程主机上以不同的操作模式启动。

• 选项 1：在调试模式下启动新应用程序。

  要在调试模式下启动新应用程序，请按如下方式调用 cuda-gdb 服务器

  $ cuda-gdbserver :1234 app_invocation
  

  其中 1234 是 TCP 端口号，cuda-gdbserver 将监听来自 cuda-gdb 的传入连接，app-invocation 是启动应用程序的调用命令，包括参数。

• 选项 2：将 cuda-gdbserver 附加到正在运行的进程

  要将 cuda-gdbserver 附加到已在运行的进程，必须使用 --attach 选项，后跟进程标识号 (PID)

  $ cuda-gdbserver :1234 --attach 5678
  

  其中 1234 是 TCP 端口号，5678 是 cuda-gdbserver 必须附加到的应用程序的进程标识符。

注意

在 QNX 平台上不支持附加到已在运行的进程。

在客户端上启动 cuda-gdb

配置 cuda-gdb 以使用以下任一方式连接到远程目标

(cuda-gdb) target remote

或

(cuda-gdb) target extended-remote

注意

QNX 平台需要将 QNX_TARGET 环境变量设置为目标 rootfs 的位置。如果这在运行 cuda-gdb 的客户端上不可用，请将其设置为空字符串。

如果调试目标上安装的库可能与调试主机上安装的库不同，建议使用 set sysroot 命令。例如，可以将 cuda-gdb 配置为按如下方式连接到远程目标

(cuda-gdb) set sysroot remote://
(cuda-gdb) target remote 192.168.0.2:1234

其中 192.168.0.2 是远程目标的 IP 地址或域名，1234 是先前由 cuda-gdbserver 打开的 TCP 端口。

3.3.4. 多个调试器

对于计算能力为 6.0 及更高的设备，可以同时进行多个调试会话。

对于计算能力低于 6.0 的设备，只要独占使用 CUDA 设备，就可以同时进行多个调试会话。例如，一个 CUDA-GDB 实例可以调试使用第一个 GPU 的第一个应用程序，而另一个 CUDA-GDB 实例可以调试使用第二个 GPU 的第二个应用程序。通过使用 CUDA_VISIBLE_DEVICES 环境变量指定哪个 GPU 对应用程序可见，即可实现 GPU 的独占使用。

$ CUDA_VISIBLE_DEVICES=1 cuda-gdb my_app

此外，对于计算能力低于 6.0 的设备，在启用软件抢占 (set cuda software_preemption on) 的情况下，可以使用多个 CUDA-GDB 实例来调试在同一 GPU 上上下文切换的 CUDA 应用程序。

3.3.5. 附加/分离

CUDA-GDB 可以使用 GDB 的内置命令附加到和分离计算能力为 2.0 及更高版本的 GPU 上运行的 CUDA 应用程序。

此外，如果在运行 CUDA 应用程序之前将环境变量 CUDA_DEVICE_WAITS_ON_EXCEPTION 设置为 1，则应用程序将正常运行，直到发生设备异常。然后应用程序将等待 CUDA-GDB 附加到它以进行进一步调试。WSL 不支持此功能。

注意

默认情况下，在某些 Linux 发行版上，由于安全设置，调试器无法附加到已在运行的进程。为了启用 CUDA 调试器的附加功能，应以 root 身份启动 cuda-gdb，或者应使用以下命令将 /proc/sys/kernel/yama/ptrace_scope 设置为零

$ sudo sh -c "echo 0 >/proc/sys/kernel/yama/ptrace_scope"

要使更改永久生效，请编辑 /etc/sysctl.d/10-ptrace.conf。

4. CUDA-GDB 扩展

4.1. 命令命名约定

现有的 GDB 命令保持不变。每个新的 CUDA 命令或选项都以 CUDA 关键字为前缀。CUDA-GDB 命令名称将尽可能类似于用于调试主机代码的等效 GDB 命令。例如，用于显示主机线程和切换到主机线程 1 的 GDB 命令分别是

(cuda-gdb) info threads
(cuda-gdb) thread 1

要显示 CUDA 线程并切换到 cuda 线程 1，用户只需键入

(cuda-gdb) info cuda threads
(cuda-gdb) cuda thread 1

4.2. 获取帮助

与 GDB 命令一样，可以使用 help 命令从 cuda-gdb 命令行访问 CUDA 命令的内置帮助

(cuda-gdb) help cuda name_of_the_cuda_command
(cuda-gdb) help set cuda name_of_the_cuda_option
(cuda-gdb) help info cuda name_of_the_info_cuda_command

此外，与任何其他 GDB 命令一样，可以通过按 TAB 键自动完成所有 CUDA 命令。

也可以使用 apropos 命令查询 CUDA 命令。

4.3. 初始化文件

CUDA-GDB 的初始化文件名为 .cuda-gdbinit，并遵循与 GDB 使用的标准 .gdbinit 文件相同的规则。初始化文件可以包含任何 CUDA-GDB 命令。当启动 CUDA-GDB 时，这些命令将按顺序处理。

4.4. GUI 集成

Emacs

CUDA-GDB 可与 Emacs 和 XEmacs 中的 GUD 一起使用。除了指向正确的二进制文件外，无需其他步骤。

要使用 CUDA-GDB，必须将 gud-gdb-command-name 变量设置为 cuda-gdb annotate=3。使用 M-x customize-variable 设置变量。

确保 cuda-gdb 存在于 Emacs/XEmacs $PATH 中。

DDD

CUDA-GDB 可与 DDD 一起使用。要将 DDD 与 CUDA-GDB 一起使用，请使用以下命令启动 DDD

ddd --debugger cuda-gdb

cuda-gdb 必须在您的 $PATH 中。

4.5. GPU 核心转储支持

有两种方法可以配置 CUDA 应用程序的核心转储选项。在应用程序环境中设置环境变量，或使用 CUDA 驱动程序 API 从应用程序以编程方式进行配置。

注意

当其他 CUDA 开发人员工具（包括 CUDA-GDB）与应用程序交互时，不支持 GPU 核心转储生成，除非明确记录为受支持的用例（例如，generate-core-file 命令）。

为 GPU 核心转储生成进行编译

无论用于生成 GPU 应用程序的编译标志如何，都会生成 GPU 核心转储。为了获得最佳调试体验，建议使用 NVCC 的 -g -G 或 -lineinfo 选项编译应用程序。有关传递调试编译标志的更多信息，请参阅 编译应用程序。

使用环境变量在异常时启用 GPU 核心转储生成

将 CUDA_ENABLE_COREDUMP_ON_EXCEPTION 环境变量设置为 1 以在遇到 GPU 异常时启用生成 GPU 核心转储。默认情况下禁用此选项。

将 CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION 环境变量设置为 0 以在遇到 GPU 异常时禁用生成 CPU 核心转储。启用 GPU 核心转储生成时，默认启用此选项，但已弃用，建议使用 CUDA_COREDUMP_GENERATION_FLAGS="skip_abort" 选项。

将 CUDA_ENABLE_LIGHTWEIGHT_COREDUMP 环境变量设置为 1 以启用生成轻量级核心文件而不是完整核心文件。启用后，GPU 核心转储将不包含应用程序的内存转储（本地、共享、全局）。默认情况下禁用此选项，并已弃用，建议使用以下 CUDA_COREDUMP_GENERATION_FLAGS 设置。

注意

即使未遇到异常，启用核心转储生成也可能会影响应用程序性能。

控制 GPU 核心转储生成的行为

可以使用 CUDA_COREDUMP_GENERATION_FLAGS 环境变量在生成 GPU core dump 时偏离默认生成行为。可以为此环境变量提供多个标志，并以 , 分隔。这些标志可用于完成诸如减小生成的 GPU core dump 大小或偏离默认值的其他所需行为等任务。下表列出了每个标志及其存在时的行为。

GPU core dump CUDA_COREDUMP_GENERATION_FLAGS

环境变量标志	描述
skip_nonrelocated_elf_images	禁用在 GPU core dump 中包含未重定位的 elf 镜像副本。仅存在重定位的镜像。
skip_global_memory	禁用转储 GPU 全局和 constbank 内存段。
skip_shared_memory	禁用转储 GPU 共享内存段。
skip_local_memory	禁用转储 GPU 本地内存段。
skip_abort	禁用在 GPU core dump 生成过程结束时调用 abort()。

注意

将 CUDA_ENABLE_LIGHTWEIGHT_COREDUMP 环境变量设置为 1 等同于 CUDA_COREDUMP_GENERATION_FLAGS="skip_nonrelocated_elf_images,skip_global_memory,skip_shared_memory,skip_local_memory"。

注意

将 CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION 环境变量设置为 0 等同于 CUDA_COREDUMP_GENERATION_FLAGS="skip_abort"。

core dump 生成的限制和注意事项

以下限制适用于 core dump 支持

• 对于 Windows WDDM，仅在计算能力为 6.0 或更高的 GPU 上支持 GPU core dump。Windows TCC 在所有受支持的计算能力上都支持 GPU core dump。

• 在适用于 Linux 的 Windows 子系统上，对于在 SLI 模式下运行的 GPU，不支持 GPU core dump。支持多 GPU 设置，但在驱动程序控制面板中无法启用 SLI 模式。

• 仅当启用 硬件调度模式 时，适用于 Linux 的 Windows 子系统才支持 GPU core dump。

• 仅当执行的二进制文件位于 Linux 文件系统上时，适用于 Linux 的 Windows 子系统才支持 GPU core dump，这是由于执行 Linux 程序时存在 NTFS 限制。

• 对于适用于 Linux 的 Windows 子系统，不支持使用 CUDA_ENABLE_USER_TRIGGERED_COREDUMP 选项生成 GPU core dump。

• 在 QNX 平台上，当前不支持使用 CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION 生成 CPU core dump。

• NVIDIA CMP 产品线不支持 GPU core dump。

• 每个上下文的 core dump 只能在计算能力为 6.0 或更高的 GPU 上启用。当使用 Coredump 属性控制 API 时，计算能力低于 6.0 的 GPU 将返回 CUDA_ERROR_NOT_SUPPORTED。

• 如果 MPS 客户端触发 core dump，则在同一 MPS 服务器上运行的所有其他客户端都将发生故障。如果间接发生故障的客户端启用了 core dump 生成，它们也将生成 core dump。

• 当其他开发人员工具（包括 CUDA-GDB）与应用程序交互时，不支持 GPU core dump。除非明确记录为受支持的用例（例如 generate-cuda-core-file 命令）。

• 当在异常时生成 core dump 时，如果内核在异常被识别之前退出，则可能导致 corefile 生成失败。有关如何解决此问题的策略，请参阅 GPU 错误报告 中的注释。

注意

用户不应向应用程序进程发送信号，并确保应用程序进程在 core dump 生成过程中不会自动终止。这样做可能会导致 GPU coredump 生成中止。

注意

从 CUDA 11.6 开始，compute-sanitizer 工具可以在检测到错误时通过使用 --generate-coredump yes 选项来生成 GPU core dump。生成 core dump 后，目标应用程序将中止。有关更多信息，请参阅 compute-sanitizer 文档：https://docs.nvda.net.cn/compute-sanitizer/ComputeSanitizer/index.html#coredump

注意

CPU core dump 将位于特定于发行版的位置。检查 /proc/sys/kernel/core_pattern 文件通常会提示 CPU core dump 的名称/位置。

注意

NVIDIA vGPU 平台必须显式启用调试支持才能执行 GPU core dump 生成。请参考 虚拟 GPU 软件用户指南，了解如何在 vGPU 上启用调试。

注意

NVIDIA Jetson 和 Drive Tegra 设备必须显式启用调试支持才能执行 GPU core dump 生成。请参阅 在 Jetson 和 Drive Tegra 设备上使用 CUDA-GDB 调试器 部分。

注意

当在运行 QNX 的 NVIDIA Drive Tegra 设备上生成 core dump 时，生成 CPU core dump 时可能会挂起。如果遇到挂起，请将 CUDA_ENABLE_CPU_COREDUMP_EXCEPTION 设置为 0。

注意

如果在运行使用 OptiX/RTCore 构建的程序时未生成 core dump，请尝试将环境变量 OPTIX_FORCE_DEPRECATED_LAUNCHER 设置为 1。请参阅 调试 OptiX/RTCore 应用程序 部分。

注意

Windows 和适用于 Linux 的 Windows 子系统 (WSL) 上的 Core dump 生成需要通过将注册表项 >HKEY_LOCAL_MACHINE\SOFTWARE\NVIDIA Corporation\GPUDebugger\EnableInterface 设置为 (DWORD) 1 来启用调试接口。

注意

仅在 devtools 模式下，才支持在以 Confidential Compute 模式运行的 GPU 上进行 GPU core dump。有关如何启用该模式的更多详细信息，请参阅 Confidential Computing Deployment Guide <https://docs.nvda.net.cn/confidential-computing-deployment-guide.pdf>。

GPU core dump 文件的命名

默认情况下，GPU core dump 在当前工作目录中创建。它被命名为 core_TIME_HOSTNAME_PID.nvcudmp，其中 TIME 是自 Epoch 以来经过的秒数，HOSTNAME 是运行 CUDA 应用程序的机器的主机名，PID 是 CUDA 应用程序的进程标识符。

可以使用 CUDA_COREDUMP_FILE 环境变量来定义用于更改 GPU core dump 文件名称的模板。该模板可以是绝对路径，也可以是相对于当前工作目录的相对路径。该模板可以包含 % 说明符，在创建 GPU core dump 时，这些说明符将替换为以下模式

说明符	描述
%h	运行 CUDA 应用程序的机器的主机名
%p	CUDA 应用程序的进程标识符
%t	时间，以自 Epoch 以来经过的秒数表示，1970-01-01 00:00:00 +0000 (UTC)

例如，将 CUDA_COREDUMP_FILE 设置为

export CUDA_COREDUMP_FILE=newName.%h.%p

将导致 GPU core dump 被写入相对于当前工作目录的 newName.myhost.1234。这里 myhost 和 1234 分别被替换为真实的主机名和 pid。

将 CUDA_COREDUMP_FILE 设置为

export CUDA_COREDUMP_FILE="/home/$USER/newName.%h.%p"

将导致 GPU core dump 被写入用户的主目录，其命名逻辑与上述示例相同。

如果 CUDA_COREDUMP_FILE 指向 FIFO 类型的现有文件（例如，命名管道），则 core dump 将被流式传输到该文件。

可以通过 CUDA_COREDUMP_FILE 将 Coredump 通过管道传输到 shell 命令，格式如下

export CUDA_COREDUMP_FILE='| cmd > file'

例如，要将 coredump 通过管道传输到 gzip，请使用

export CUDA_COREDUMP_FILE='| gzip -9 > cuda-coredump.gz'

注意

当通过管道传输 coredump 时，% 说明符将无法识别。

启用用户诱导的 GPU core dump 生成

对于支持计算抢占的设备，用户可以中断正在运行的 CUDA 进程以生成 GPU core dump。

将 CUDA_ENABLE_USER_TRIGGERED_COREDUMP 环境变量设置为 1，以便启用生成用户诱导的 GPU core dump。默认情况下，此选项处于禁用状态。设置此环境变量将为每个后续运行的 CUDA 进程打开一个通信管道。要诱导 GPU core dump，用户只需写入管道即可。

要更改默认管道文件名，请将 CUDA_COREDUMP_PIPE 环境变量设置为特定的管道名称。默认管道名称的格式如下：corepipe.cuda.HOSTNAME.PID，其中 HOSTNAME 是运行 CUDA 应用程序的机器的主机名，PID 是 CUDA 应用程序的进程标识符。此环境变量可以采用上一节中描述的 % 说明符。

可以使用 cuda-gdb generate-core-file 或 gcore 命令手动生成 GPU core dump。可以传递可选的文件名，有关详细信息，请参阅 help generate-core-file。有多个标志可用于控制要保存的内容以及是否同时生成 CPU core 文件。

显示 core dump 生成进度

默认情况下，当应用程序崩溃并生成 GPU core dump 时，应用程序可能会显得无响应或冻结，直到完全生成为止。

将 CUDA_COREDUMP_SHOW_PROGRESS 环境变量设置为 1，以便将 core dump 生成进度消息打印到 stderr。这可以用于确定 coredump 生成的进度。

coredump: SM 1/14 has finished state collection
coredump: SM 2/14 has finished state collection
coredump: SM 3/14 has finished state collection
coredump: SM 4/14 has finished state collection
coredump: SM 5/14 has finished state collection
coredump: SM 6/14 has finished state collection
coredump: SM 7/14 has finished state collection
coredump: SM 8/14 has finished state collection
coredump: SM 9/14 has finished state collection
coredump: SM 10/14 has finished state collection
coredump: SM 11/14 has finished state collection
coredump: SM 12/14 has finished state collection
coredump: SM 13/14 has finished state collection
coredump: SM 14/14 has finished state collection
coredump: Device 1/1 has finished state collection
coredump: Calculating ELF file layout
coredump: ELF file layout calculated
coredump: Writing ELF file to core_TIME_HOSTNAME_PID.nvcudmp
coredump: Writing out global memory (1073741824 bytes)
coredump: 5%...
coredump: 10%...
coredump: 15%...
coredump: 20%...
coredump: 25%...
coredump: 30%...
coredump: 35%...
coredump: 40%...
coredump: 45%...
coredump: 50%...
coredump: 55%...
coredump: 60%...
coredump: 65%...
coredump: 70%...
coredump: 75%...
coredump: 80%...
coredump: 85%...
coredump: 90%...
coredump: 95%...
coredump: 100%...
coredump: Writing out device table
coredump: Finalizing
coredump: All done

使用 CUDA Driver API 启用 GPU core dump 生成

Driver API 具有所有环境变量的等效设置，并且增加了能够为每个上下文而不是全局设置不同 core dump 设置的功能。此 API 可以直接在您的应用程序内部调用。使用 cuCoredumpGetAttributeGlobal 和 cuCoredumpSetAttributeGlobal 来获取或设置全局属性。使用 cuCoredumpGetAttribute 和 cuCoredumpSetAttribute 来获取或设置每个上下文的属性。有关更多信息，请参阅 Coredump 属性控制 API 手册。

下表列出了环境变量和等效的 CUcoredumpSettings 标志，这些标志可用于通过 Coredump 属性控制 API 管理 core dump 设置。

注意

CU_COREDUMP_ENABLE_USER_TRIGGER 设置只能在驱动程序 API 中全局设置，并且必须在启用用户触发的 core dump 之前设置 CU_COREDUMP_PIPE（如果需要）。

GPU core dump 配置参数

环境变量	描述
环境变量 CUDA_ENABLE_COREDUMP_ON_EXCEPTION CUcoredumpSettings 标志 CU_COREDUMP_ENABLE_ON_EXCEPTION	为异常启用 GPU core dump 生成。默认情况下禁用。
环境变量 CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION CUcoredumpSettings 标志 CU_COREDUMP_TRIGGER_HOST	在 GPU core dump 完成后触发主机 (CPU) core dump。默认情况下启用。自 CUDA 12.5 起已弃用，转而支持与 CU_COREDUMP_GENERATION_FLAGS 选项一起给出的 CU_COREDUMP_SKIP_ABORT 值。
环境变量 CUDA_ENABLE_LIGHTWEIGHT_COREDUMP CUcoredumpSettings 标志 CU_COREDUMP_LIGHTWEIGHT	启用后，GPU core dump 将不包含应用程序的内存转储（本地、共享、全局）。默认情况下禁用。自 CUDA 12.5 起已弃用，转而支持与 CU_COREDUMP_GENERATION_FLAGS 选项一起给出的 CU_COREDUMP_LIGHTWEIGHT_FLAGS 值。
环境变量 CUDA_ENABLE_USER_TRIGGERED_COREDUMP CUcoredumpSettings 标志 CU_COREDUMP_ENABLE_USER_TRIGGER	通过写入在 COREDUMP_PIPE 设置中定义的管道来启用用户可触发的 core dump。默认情况下禁用。
环境变量 CUDA_COREDUMP_FILE CUcoredumpSettings 标志 CU_COREDUMP_FILE	GPU core dump 的文件名模板。
环境变量 CUDA_COREDUMP_PIPE CUcoredumpSettings 标志 CU_COREDUMP_PIPE	用户管道触发器的文件名模板。
环境变量 CUDA_COREDUMP_GENERATION_FLAGS CUcoredumpSettings 标志 CU_COREDUMP_GENERATION_FLAGS	用于控制 GPU coredump 生成设置的标志。使用按位 OR 在一次调用中传递多个设置。有关当前版本中有效值的列表，请参阅 cuda.h 中的 CUCoredumpGenerationFlags。

在 cuda-gdb 中检查 GPU 和 GPU+CPU core dump

使用以下命令将 GPU core dump 加载到调试器中

• (cuda-gdb) target cudacore core.cuda.localhost.1234
  

  这将打开 core dump 文件并打印程序执行期间遇到的异常。然后，发出标准 cuda-gdb 命令以进一步调查中止时设备上的应用程序状态。

使用以下命令将 CPU 和 GPU core dump 加载到调试器中

• (cuda-gdb) target core core.cpu core.cuda
  

  这将打开 core dump 文件并打印程序执行期间遇到的异常。然后，发出标准 cuda-gdb 命令以进一步调查中止时主机和设备上的应用程序状态。

注意

Coredump 检查不需要系统上安装 GPU

5. 内核焦点

一个 CUDA 应用程序可能正在运行多个主机线程和许多设备线程。为了简化应用程序状态信息的可视化，命令将应用于焦点中的实体。

当焦点设置为主机线程时，命令将仅应用于该主机线程（除非应用程序完全恢复，例如）。在设备端，焦点始终设置为最低粒度级别——设备线程。

5.1. 软件坐标与硬件坐标

设备线程属于一个块，该块又属于一个内核。线程、块和内核是焦点的软件坐标。设备线程在 lane 上运行。lane 属于 warp，warp 属于 SM，SM 又属于设备。Lane、warp、SM 和设备是焦点的硬件坐标。只要它们保持一致，软件坐标和硬件坐标可以互换和同时使用。

有时会使用另一个软件坐标：grid。grid 和内核之间的区别在于范围。grid ID 在每个 GPU 上是唯一的，而内核 ID 在所有 GPU 上是唯一的。因此，内核与（grid，device）元组之间存在 1:1 映射。

注意

如果启用了软件抢占（set cuda software_preemption on），则设备线程对应的硬件坐标可能会在设备上恢复执行后发生更改。但是，软件坐标将保持不变，并且在设备线程的生命周期内不会更改。

5.2. 当前焦点

要检查当前焦点，请使用 cuda 命令，后跟感兴趣的坐标

(cuda-gdb) cuda device sm warp lane block thread
block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0
(cuda-gdb) cuda kernel block thread
kernel 1, block (0,0,0), thread (0,0,0)
(cuda-gdb) cuda kernel
kernel 1

5.3. 切换焦点

要切换当前焦点，请使用 cuda 命令，后跟要更改的坐标

(cuda-gdb) cuda device 0 sm 1 warp 2 lane 3
[Switching focus to CUDA kernel 1, grid 2, block (8,0,0), thread
(67,0,0), device 0, sm 1, warp 2, lane 3]
374 int totalThreads = gridDim.x * blockDim.x;

如果命令未完全定义指定的焦点，则调试器将假定省略的坐标设置为当前焦点中的坐标，包括块和线程的子坐标。

(cuda-gdb) cuda thread (15)
[Switching focus to CUDA kernel 1, grid 2, block (8,0,0), thread
(15,0,0), device 0, sm 1, warp 0, lane 15]
374 int totalThreads = gridDim.x * blockDim.x;

块和线程参数的括号是可选的。

(cuda-gdb) cuda block 1 thread 3
[Switching focus to CUDA kernel 1, grid 2, block (1,0,0), thread (3,0,0),
device 0, sm 3, warp 0, lane 3]
374 int totalThreads = gridDim.x * blockDim.

6. 程序执行

在 CUDA-GDB 中，应用程序的启动方式与 GDB 相同，都是使用 run 命令。本章介绍如何中断和单步执行 CUDA 应用程序

6.1. 中断应用程序

如果 CUDA 应用程序似乎挂起或卡在无限循环中，则可以通过按 CTRL+C 手动中断应用程序。当收到信号时，GPU 将被暂停，并且将出现 cuda-gdb 提示符。

此时，可以由用户自行决定检查、修改、单步执行、恢复或终止程序。

此功能仅限于在调试器内运行的应用程序。无法中断和调试在调试器外部启动的应用程序。

6.2. 单步执行

支持单步执行设备代码。但是，与主机代码单步执行不同，设备代码单步执行在 warp 级别工作。这意味着单步执行设备内核会推进当前焦点中 warp 中的所有活动线程。warp 中的发散线程不会单步执行。当焦点中的 CUDA 线程变为发散时，行为取决于 set cuda step_divergent_lanes 的值。当设置为 on（默认值）时，将连续单步执行焦点中的 warp，直到焦点中的 CUDA 线程变为活动状态。当设置为 off 时，将单步执行焦点中的 warp，并将焦点 CUDA 线程更改为 warp 中最近的活动 lane。

为了推进多个 warp 的执行，必须在所需位置设置断点，然后必须完全恢复应用程序。

一个特殊情况是单步执行线程屏障调用，例如：__syncthreads() 或集群范围的屏障。在这种情况下，将在屏障之后立即设置隐式临时断点，并且所有线程都将恢复，直到命中临时断点。

只要设备函数不是内联的，您就可以步入、步过或步出设备函数。要强制编译器不内联函数，必须将 __noinline__ 关键字添加到函数声明中。

无法保证在先前 PC 上执行的设备上的异步 SASS 指令已完成。

使用动态并行性，可以直接从设备代码调用多个 CUDA API。以下列表定义了遇到这些 API 时的单步执行行为

• 当遇到设备端内核启动（由 <<<>>> 启动语法表示）时，step 和 next 命令将具有相同的行为，并且都将步过启动调用。

• 在 Hopper 之前的设备（SM 9.0）上，步入已弃用的 cudaDeviceSynchronize() 会导致未定义的行为。用户应改为步过此调用。

• 当单步执行设备 grid 启动完成时，焦点将自动切换回 CPU。cuda kernel 焦点切换命令必须用于切换到另一个感兴趣的 grid（如果仍然存在）。

注意

无法步入设备启动调用（以及调用启动的例程）。

7. 断点和观察点

有多种方法可以在 CUDA 应用程序上设置断点。以下描述了这些方法。用于在设备代码上设置断点的命令与用于在主机代码上设置断点的命令相同。

如果在设备代码上设置了断点，则该断点将被标记为挂起，直到内核的 ELF 镜像加载完成。此时，断点将被解析，其地址将被更新。

设置断点后，它会强制所有驻留的 GPU 线程在到达相应的 PC 时在此位置停止。

当一个线程命中断点时，不能保证其他线程会同时命中该断点。因此，同一个断点可能会被命中多次，用户必须小心检查哪些线程实际命中该断点。disable 命令可用于防止其他线程命中该断点。

7.1. 符号断点

要在函数入口处设置断点，请使用 break 命令，后跟函数或方法的名称

(cuda-gdb) break my_function
(cuda-gdb) break my_class::my_method

对于模板化函数和方法，必须给出完整的签名

(cuda-gdb) break int my_templatized_function<int>(int)

也可以使用函数的 mangled name。要查找函数的 mangled name，可以使用以下命令

(cuda-gdb) set demangle-style none
(cuda-gdb) info function my_function_name
(cuda-gdb) set demangle-style auto

7.2. 行断点

要在特定行号上设置断点，请使用以下语法

(cuda-gdb) break my_file.cu:185

如果指定的行对应于模板化代码中的指令，则将创建多个断点，每个模板化代码实例一个断点。

7.3. 地址断点

要在特定地址设置断点，请使用 break 命令，并将地址作为参数

(cuda-gdb) break *0x1afe34d0

地址可以是设备或主机上的任何地址。

7.4. 内核入口断点

要在每个启动的内核的第一条指令处中断，请将 break_on_launch 选项设置为 application

(cuda-gdb) set cuda break_on_launch application

有关更多信息，请参阅 set cuda break_on_launch。

7.5. 条件断点

要使断点成为条件断点，请使用可选的 if 关键字或 cond 命令。

(cuda-gdb) break foo.cu:23 if threadIdx.x == 1 && i < 5
(cuda-gdb) cond 3 threadIdx.x == 1 && i < 5

条件表达式可以引用任何变量，包括内置变量，例如 threadIdx 和 blockIdx。条件表达式中不允许函数调用。

请注意，条件断点始终被命中和评估，但仅当条件语句评估为 TRUE 时，调试器才报告断点被命中。命中断点并评估相应条件语句的过程非常耗时。因此，在使用条件断点时运行应用程序可能会减慢调试会话的速度。此外，如果条件语句始终评估为 FALSE，则调试器可能显得挂起或卡住，但事实并非如此。您可以使用 CTRL-C 中断应用程序以验证是否正在进行进度。

可以在尚未加载的 CUDA 模块的代码上设置条件断点。然后，仅当加载该模块的 ELF 镜像时才会进行条件验证。因此，条件表达式中的任何错误都将延迟到加载 CUDA 模块时。要仔细检查所需的条件表达式，请首先在所需位置设置无条件断点并继续。当命中断点时，使用 cond 命令评估所需的条件语句。

7.6. 观察点

不支持 CUDA 代码上的观察点。

支持主机代码上的观察点。建议用户阅读 GDB 文档，以获取有关如何在主机代码上设置观察点的教程。

8. 检查程序状态

8.1. 内存和变量

GDB print 命令已扩展为可以解析任何程序变量的位置，并且可以用于显示任何 CUDA 程序变量的内容，包括

• 通过 cudaMalloc() 分配的数据

• 驻留在各种 GPU 内存区域（例如，共享内存、本地内存和全局内存）中的数据

• 特殊的 CUDA 运行时变量，例如 threadIdx

8.2. 变量存储和可访问性

根据变量类型和用法，变量可以存储在寄存器中，也可以存储在 local、shared、const 或 global 内存中。您可以打印任何变量的地址，以找出其存储位置并直接访问关联的内存。

下面的示例显示了如何直接访问类型为 shared int * 的变量数组，以便查看数组中存储的值。

(cuda-gdb) print &array
$1 = (@shared int (*)[0]) 0x20
(cuda-gdb) print array[0]@4
$2 = {0, 128, 64, 192}

您还可以访问索引到起始偏移量的共享内存，以查看存储的值

(cuda-gdb) print *(@shared int*)0x20
$3 = 0
(cuda-gdb) print *(@shared int*)0x24
$4 = 128
(cuda-gdb) print *(@shared int*)0x28
$5 = 64

下面的示例显示了如何访问内核输入参数的起始地址。

(cuda-gdb) print &data
$6 = (const @global void * const @parameter *) 0x10
(cuda-gdb) print *(@global void * const @parameter *) 0x10
$7 = (@global void * const @parameter) 0x110000</>

8.3. Info CUDA 命令

这些命令显示有关 GPU 和应用程序 CUDA 状态的信息。可用选项包括

devices

有关所有设备的信息

sms

有关当前设备中所有活动 SM 的信息

warps

有关当前 SM 中所有活动 warp 的信息

lanes

有关当前 warp 中所有活动 lane 的信息

kernels

有关所有活动内核的信息

blocks

有关当前内核中所有活动块的信息

threads

有关当前内核中所有活动线程的信息

launch trace

有关焦点内核的父内核的信息

launch children

有关焦点内核启动的内核的信息

contexts

有关所有上下文的信息

过滤器可以应用于每个 info cuda 命令。过滤器限制命令的范围。过滤器由一个或多个限制组成。限制可以是以下任何一种

• device n

• sm n

• warp n

• lane n

• kernel n

• grid n

• block x[,y] 或 block (x[,y])

• thread x[,y[,z]] 或 thread (x[,y[,z]])

• breakpoint all 和 breakpoint n

其中 n、x、y、z 是整数，或以下特殊关键字之一：current、any 和 all。current 表示应使用当前焦点中的相应值。any 和 all 表示任何值都是可接受的。

注意

breakpoint all 和 breakpoint n 过滤器仅对 info cuda threads 命令有效。

8.3.1. info cuda devices

此命令枚举系统中按设备索引排序的所有 GPU。* 表示当前焦点中的设备。此命令支持过滤器。默认值为 device all。如果未找到任何活动的 GPU，则此命令会打印 No CUDA Devices。在遇到第一个内核启动之前，设备不被视为活动状态。

(cuda-gdb) info cuda devices
  Dev PCI Bus/Dev ID                Name Description SM Type SMs Warps/SM Lanes/Warp Max Regs/Lane Active SMs Mask
    0        06:00.0 GeForce GTX TITAN Z      GK110B   sm_35  15       64         32           256 0x00000000
    1        07:00.0 GeForce GTX TITAN Z      GK110B   sm_35  15       64         32           256 0x00000000

8.3.2. info cuda sms

此命令显示设备的全部 SM 以及 SM 上的关联活动 warp。此命令支持过滤器，默认值为 device current sm all。* 表示焦点 SM。结果按设备分组。

(cuda-gdb) info cuda sms
 SM Active Warps Mask
Device 0
* 0 0xffffffffffffffff
  1 0xffffffffffffffff
  2 0xffffffffffffffff
  3 0xffffffffffffffff
  4 0xffffffffffffffff
  5 0xffffffffffffffff
  6 0xffffffffffffffff
  7 0xffffffffffffffff
  8 0xffffffffffffffff
...

8.3.3. info cuda warps

此命令使您更深入一层，并打印焦点 SM 的所有 warp 信息。此命令支持过滤器，默认值为 device current sm current warp all。该命令可用于显示哪个 warp 执行哪个块。

(cuda-gdb) info cuda warps
Wp /Active Lanes Mask/ Divergent Lanes Mask/Active Physical PC/Kernel/BlockIdx
Device 0 SM 0
* 0    0xffffffff    0x00000000 0x000000000000001c    0    (0,0,0)
  1    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)
  2    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)
  3    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)
  4    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)
  5    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)
  6    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)
  7    0xffffffff    0x00000000 0x0000000000000000    0    (0,0,0)
 ...

8.3.4. info cuda lanes

此命令显示焦点 Warp 中的所有 Lane（线程）。此命令支持过滤器，默认过滤器为 device current sm current warp current lane all。在下面的示例中，您可以看到所有 Lane 都位于相同的物理 PC 上。此命令可用于显示哪个 Lane 执行哪个线程。

(cuda-gdb) info cuda lanes
  Ln    State  Physical PC        ThreadIdx
Device 0 SM 0 Warp 0
*  0    active 0x000000000000008c   (0,0,0)
   1    active 0x000000000000008c   (1,0,0)
   2    active 0x000000000000008c   (2,0,0)
   3    active 0x000000000000008c   (3,0,0)
   4    active 0x000000000000008c   (4,0,0)
   5    active 0x000000000000008c   (5,0,0)
   6    active 0x000000000000008c   (6,0,0)
   7    active 0x000000000000008c   (7,0,0)
   8    active 0x000000000000008c   (8,0,0)
   9    active 0x000000000000008c   (9,0,0)
  10    active 0x000000000000008c  (10,0,0)
  11    active 0x000000000000008c  (11,0,0)
  12    active 0x000000000000008c  (12,0,0)
  13    active 0x000000000000008c  (13,0,0)
  14    active 0x000000000000008c  (14,0,0)
  15    active 0x000000000000008c  (15,0,0)
  16    active 0x000000000000008c  (16,0,0)
 ...

8.3.5. info cuda kernels

此命令在焦点 GPU 上显示所有活动的 Kernel。它打印每个 Kernel 的 SM 掩码、Kernel ID 和 Grid ID，以及相关的维度和参数。“Parent”列显示父 Grid 的 Kernel ID。此命令支持过滤器，默认过滤器为 kernel all。

(cuda-gdb) info cuda kernels
  Kernel Parent Dev Grid Status   SMs Mask   GridDim  BlockDim      Name Args
*      1      -   0    2 Active 0x00ffffff (240,1,1) (128,1,1) acos_main parms=...

此命令还将显示已在 GPU 上通过动态并行性启动的 Grid。Grid ID 为负值的 Kernel 是从 GPU 启动的，而 Grid ID 为正值的 Kernel 是从 CPU 启动的。

8.3.6. info cuda blocks

此命令显示焦点 Kernel 的所有活动或正在运行的 Block。结果按 Kernel 分组。此命令支持过滤器，默认过滤器为 kernel current block all。默认情况下，输出是合并的。

(cuda-gdb) info cuda blocks
   BlockIdx   To BlockIdx  Count  State
Kernel 1
*  (0,0,0)    (191,0,0)    192    running

可以按如下方式关闭合并，在这种情况下，将显示有关设备和 SM 的更多信息

(cuda-gdb) set cuda coalescing off

以下是关闭合并时同一命令的输出。

(cuda-gdb) info cuda blocks
  BlockIdx   State    Dev SM
Kernel 1
*   (0,0,0)   running   0   0
    (1,0,0)   running   0   3
    (2,0,0)   running   0   6
    (3,0,0)   running   0   9
    (4,0,0)   running   0  12
    (5,0,0)   running   0  15
    (6,0,0)   running   0  18
    (7,0,0)   running   0  21
    (8,0,0)   running   0   1
 ...

8.3.7. info cuda threads

此命令显示应用程序的活动 CUDA Block 和线程，以及这些 Block 中的线程总数。还显示了虚拟 PC 以及相关的源文件和行号信息。结果按 Kernel 分组。该命令支持过滤器，默认过滤器为 kernel current block all thread all。默认情况下，输出是合并的，如下所示

(cuda-gdb) info cuda threads
  BlockIdx ThreadIdx To BlockIdx ThreadIdx Count   Virtual PC    Filename   Line
Device 0 SM 0
* (0,0,0  (0,0,0)    (0,0,0)  (31,0,0)    32  0x000000000088f88c   acos.cu   376
  (0,0,0)(32,0,0)  (191,0,0) (127,0,0) 24544  0x000000000088f800   acos.cu   374
 ...

可以按如下方式关闭合并，在这种情况下，将显示更多信息和输出。

(cuda-gdb) info cuda threads
   BlockIdx  ThreadIdx  Virtual PC         Dev SM Wp Ln   Filename  Line
Kernel 1
*  (0,0,0)    (0,0,0)  0x000000000088f88c   0  0  0  0    acos.cu    376
   (0,0,0)    (1,0,0)  0x000000000088f88c   0  0  0  1    acos.cu    376
   (0,0,0)    (2,0,0)  0x000000000088f88c   0  0  0  2    acos.cu    376
   (0,0,0)    (3,0,0)  0x000000000088f88c   0  0  0  3    acos.cu    376
   (0,0,0)    (4,0,0)  0x000000000088f88c   0  0  0  4    acos.cu    376
   (0,0,0)    (5,0,0)  0x000000000088f88c   0  0  0  5    acos.cu    376
   (0,0,0)    (6,0,0)  0x000000000088f88c   0  0  0  6    acos.cu    376
   (0,0,0)    (7,0,0)  0x000000000088f88c   0  0  0  7    acos.cu    376
   (0,0,0)    (8,0,0)  0x000000000088f88c   0  0  0  8    acos.cu    376
   (0,0,0)    (9,0,0)  0x000000000088f88c   0  0  0  9    acos.cu    376
 ...

注意

在合并形式中，线程必须是连续的才能被合并。如果某些线程当前未在硬件上运行，它们将在线程范围内创建空洞。例如，如果一个 Kernel 由 2 个 Block 组成，每个 Block 有 16 个线程，并且只有最低的 8 个线程处于活动状态，则将打印 2 个合并的范围：一个范围用于 Block 0 线程 0 到 7，另一个范围用于 Block 1 线程 0 到 7。由于 Block 0 中的线程 8-15 未运行，因此无法合并这两个范围。

该命令还支持 breakpoint all 和 breakpoint breakpoint_number 作为过滤器。前者显示命中用户设置的所有 CUDA 断点的线程。后者显示命中 CUDA 断点 *breakpoint_number* 的线程。

(cuda-gdb) info cuda threads breakpoint all
  BlockIdx ThreadIdx         Virtual PC Dev SM Wp Ln        Filename  Line
Kernel 0
   (1,0,0)   (0,0,0) 0x0000000000948e58   0 11  0  0 infoCommands.cu    12
   (1,0,0)   (1,0,0) 0x0000000000948e58   0 11  0  1 infoCommands.cu    12
   (1,0,0)   (2,0,0) 0x0000000000948e58   0 11  0  2 infoCommands.cu    12
   (1,0,0)   (3,0,0) 0x0000000000948e58   0 11  0  3 infoCommands.cu    12
   (1,0,0)   (4,0,0) 0x0000000000948e58   0 11  0  4 infoCommands.cu    12
   (1,0,0)   (5,0,0) 0x0000000000948e58   0 11  0  5 infoCommands.cu    12

(cuda-gdb) info cuda threads breakpoint 2 lane 1
  BlockIdx ThreadIdx         Virtual PC Dev SM Wp Ln        Filename  Line
Kernel 0
   (1,0,0)   (1,0,0) 0x0000000000948e58   0 11  0  1 infoCommands.cu    12

8.3.8. info cuda launch trace

此命令显示焦点 Kernel 的 Kernel 启动跟踪。跟踪中的第一个元素是焦点 Kernel。下一个元素是启动此 Kernel 的 Kernel。跟踪持续到没有父 Kernel 为止。在这种情况下，Kernel 是由 CPU 启动的。

对于跟踪中的每个 Kernel，该命令打印 Kernel 在跟踪中的级别、Kernel ID、设备 ID、Grid ID、状态、Kernel 维度、Kernel 名称和 Kernel 参数。

(cuda-gdb) info cuda launch trace
  Lvl Kernel Dev Grid     Status   GridDim  BlockDim Invocation
*   0      3   0   -7     Active  (32,1,1)  (16,1,1) kernel3(c=5)
    1      2   0   -5 Terminated (240,1,1) (128,1,1) kernel2(b=3)
    2      1   0    2     Active (240,1,1) (128,1,1) kernel1(a=1)

已启动但未在 GPU 上运行的 Kernel 将具有 Pending 状态。当前在 GPU 上运行的 Kernel 将标记为 Active。等待再次变为活动的 Kernel 将显示为 Sleeping。当 Kernel 终止时，它将被标记为 Terminated。在极少数情况下，当调试器无法确定 Kernel 是挂起还是已终止时，状态将设置为 Undetermined。

此命令支持过滤器，默认过滤器为 kernel all。

注意

使用 set cuda software_preemption on，将不会报告任何 Kernel 为活动状态。

8.3.9. info cuda launch children

此命令显示焦点 Kernel 启动的非终止 Kernel 列表。对于每个 Kernel，都会显示 Kernel ID、设备 ID、Grid ID、Kernel 维度、Kernel 名称和 Kernel 参数。

(cuda-gdb) info cuda launch children
  Kernel Dev Grid GridDim BlockDim Invocation
*      3   0   -7 (1,1,1)  (1,1,1) kernel5(a=3)
      18   0   -8 (1,1,1) (32,1,1) kernel4(b=5)

此命令支持过滤器，默认过滤器为 kernel all。

8.3.10. info cuda contexts

此命令枚举所有 GPU 上运行的所有 CUDA Context。“*”表示当前处于焦点的 Context。此命令显示 Context 当前是否在设备上处于活动状态。

(cuda-gdb) info cuda contexts
     Context Dev    State
  0x080b9518   0 inactive
* 0x08067948   0   active

8.3.11. info cuda managed

此命令根据焦点显示设备上或主机上的所有静态托管变量。

(cuda-gdb) info cuda managed
Static managed variables on device 0 are:
managed_var = 3
managed_consts = {one = 1, e = 2.71000004, pi = 3.1400000000000001}

8.4. 反汇编

可以使用标准 GDB 反汇编指令（例如 x/i 和 display/i）反汇编设备 SASS 代码。

(cuda-gdb) x/4i $pc-32
   0xa689a8 <acos_main(acosParams)+824>: MOV R0, c[0x0][0x34]
   0xa689b8 <acos_main(acosParams)+840>: MOV R3, c[0x0][0x28]
   0xa689c0 <acos_main(acosParams)+848>: IMUL R2, R0, R3
=> 0xa689c8 <acos_main(acosParams)+856>: MOV R0, c[0x0][0x28]

注意

为了使反汇编指令正常工作，必须安装 cuobjdump 并存在于您的 $PATH 中。

在反汇编视图中，当前 PC 以 => 为前缀。对于 Maxwell (SM 5.0) 及更高版本的架构，如果指令触发异常，则将以 *> 为前缀。如果 PC 和 errorpc 是同一条指令，则将以 *=> 为前缀。

例如，考虑以下异常

CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x555555c08620 (memexceptions_kernel.cu:17)

Thread 1 "memexceptions" received signal CUDA_EXCEPTION_14, Warp Illegal Address.
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
0x0000555555c08fb0 in exception_kernel<<<(1,1,1),(1,1,1)>>> (data=0x7fffccc00000, exception=MMU_FAULT) at memexceptions_kernel.cu:50
50  }
(cuda-gdb)

可以使用 disas 命令查看触发异常的 PC 和 error PC。

(cuda-gdb) disas $pc,+16
Dump of assembler code from 0x555555c08fb0 to 0x555555c08fc0:
=> 0x0000555555c08fb0 <_Z16exception_kernelPv11exception_t+3504>:  ERRBAR
End of assembler dump.

(cuda-gdb) disas $errorpc,+16
Dump of assembler code from 0x555555c08620 to 0x555555c08630:
*> 0x0000555555c08620 <_Z16exception_kernelPv11exception_t+1056>:  ST.E.U8.STRONG.SYS [R6.64], R5
End of assembler dump.

8.5. 寄存器

可以使用标准 GDB 命令（例如 info registers）检查/修改设备寄存器代码。

(cuda-gdb) info registers $R0 $R1 $R2 $R3
R0             0xf0 240
R1             0xfffc48 16776264
R2             0x7800   30720
R3             0x80 128

寄存器也可以作为 $R<regnum> 内置变量访问，例如

(cuda-gdb) printf "%d %d\n", $R0*$R3, $R2
30720 30720

可以通过打印系统寄存器组或使用各自的伪名称来检查谓词和 CC 寄存器的值：$P0..$P6 和 $CC。

(cuda-gdb) info registers system
P0             0x1  1
P1             0x1  1
P2             0x0  0
P3             0x0  0
P4             0x0  0
P5             0x0  0
P6             0x1  1
CC             0x0  0

8.6. 常量 Bank

在 GPU 内存的常量地址空间中分配的内存驻留在称为常量 Bank 的二维数组中。常量 Bank 记为 c[X][Y]，其中 X 是 Bank 编号，Y 是偏移量。给定 Bank/偏移量对的内存地址通过便捷函数 $_cuda_const_bank(bank, offset) 获得。

(cuda-gdb) disass $pc,+16
Dump of assembler code from 0x7fffd5043d40 to 0x7fffd5043d50:
=> 0x00007fffd5043d40 <_Z9acos_main10acosParams+1856>:  MOV R0, c[0x0][0xc]
End of assembler dump.
(cuda-gdb) p *$_cuda_const_bank(0x0,0xc)
$1 = 8

9. 事件通知

当应用程序向前推进时，CUDA-GDB 会通知用户有关 Kernel 事件和 Context 事件。在 CUDA-GDB 中，Kernel 指的是在 GPU 上执行的设备代码，而 Context 指的是 GPU 上 Kernel 的虚拟地址空间。您可以启用 CUDA Context 和 Kernel 事件的输出，以查看活动 Context 和 Kernel 的流程。默认情况下，仅显示 Context 事件消息。

9.1. Context 事件

每当应用程序创建、推送、弹出或销毁 CUDA Context 时，CUDA-GDB 都可以选择显示通知消息。该消息包括 Context ID 和 Context 所属的设备 ID。

[Context Create of context 0xad2fe60 on Device 0]
[Context Destroy of context 0xad2fe60 on Device 0]

默认情况下，Context 事件通知已禁用。Context 事件通知策略由 context_events 选项控制。

• (cuda-gdb) set cuda context_events off
  

  CUDA-GDB 不显示 Context 事件通知消息（默认）。

• (cuda-gdb) set cuda context_events on
  

  CUDA-GDB 显示 Context 事件通知消息。

9.2. Kernel 事件

每当 CUDA-GDB 获知 CUDA Kernel 的启动或终止时，都可以显示通知消息。该消息包括 Kernel ID、Kernel 名称和 Kernel 所属的设备。

[Launch of CUDA Kernel 1 (kernel3) on Device 0]
[Termination of CUDA Kernel 1 (kernel3) on Device 0]

Kernel 事件通知策略由 kernel_events 和 kernel_events_depth 选项控制。

• (cuda-gdb) set cuda kernel_events none
  

可能的选项包括

none

无 Kernel，应用程序或系统（默认）

application

用户应用程序启动的 Kernel

system

驱动程序启动的任何 Kernel，例如 memset

all

任何 Kernel、应用程序和系统

• (cuda-gdb) set cuda kernel_events_depth 0
  

  控制 Kernel 的最大深度，超过该深度后将不再显示 Kernel 事件通知。值为零表示没有最大值，并且显示所有 Kernel 通知。值为一表示调试器将仅显示从 CPU 启动的 Kernel 的 Kernel 事件通知（默认）。

10. 自动错误检查

10.1. 检查 API 错误

CUDA-GDB 可以自动检查任何驱动程序 API 或运行时 API 调用的返回值。如果返回值指示错误，调试器将停止或警告用户。

行为由 set cuda api_failures 选项控制。支持三种模式

• hide 不报告 CUDA API 调用失败

• ignore 为每个致命的 CUDA API 调用失败打印警告消息（默认）

• stop 当 CUDA API 调用返回致命错误时，应用程序停止

• ignore_all 为每个 CUDA API 调用失败打印警告消息

• stop_all 当 CUDA API 调用返回任何错误时，应用程序停止

注意

成功返回值和其他非错误返回值将被忽略。对于驱动程序 API，这些是：CUDA_SUCCESS 和 CUDA_ERROR_NOT_READY。对于运行时 API，它们是 cudaSuccess 和 cudaErrorNotReady。

10.2. GPU 错误报告

借助 CUDA-GDB 中改进的 GPU 错误报告，现在可以更轻松地识别和修复应用程序错误。下表显示了在计算能力为 sm_20 及更高版本的 GPU 上报告的新错误。

注意

在发现这些错误后继续执行应用程序可能会导致应用程序终止或不确定的结果。

注意

Warp 错误可能会导致指令在异常被识别和报告之前继续执行。报告的 $errorpc 应包含导致异常的指令的精确地址。如果 Warp 在导致异常的指令执行后退出，但在异常被识别和报告之前退出，则可能导致异常未被报告。CUDA-GDB 依赖于设备上存在的活动 Warp 才能报告异常。为了帮助避免这种未报告异常的情况

• 对于 Volta+ 架构，请使用 -G 编译应用程序。有关更多信息，请参见 编译应用程序。

• 在 Kernel 退出前添加 while(1);。这将确保异常被识别和报告。

• 依靠 compute-sanitizer memcheck 工具来捕获可能导致异常的访问。

CUDA 异常代码

异常代码	错误的精度	错误的范围	描述
CUDA_EXCEPTION_0 : “设备未知异常”	未知	GPU 上的全局错误	这是由应用程序引起的全局 GPU 错误，与下面列出的任何错误代码都不匹配。这种情况应该很少发生。可能，这可能是由于 Device Hardware Stack 溢出或 Kernel 生成异常非常接近其终止。
CUDA_EXCEPTION_1 : “已弃用”	已弃用	已弃用	此异常已弃用，应视为 CUDA_EXCEPTION_0。
CUDA_EXCEPTION_2 : “Lane 用户堆栈溢出”	精确	每个 Lane/线程错误	当线程超出其堆栈内存限制时，会发生这种情况。
CUDA_EXCEPTION_3 : “设备硬件堆栈溢出”	精确	GPU 上的全局错误	当应用程序触发全局硬件堆栈溢出时，会发生这种情况。此错误的主要原因是存在函数调用时的大量发散。
CUDA_EXCEPTION_4 : “Warp 非法指令”	精确	Warp 错误	当 Warp 中的任何线程执行了非法指令时，会发生这种情况。
CUDA_EXCEPTION_5 : “Warp 超出范围地址”	精确	Warp 错误	当 Warp 中的任何线程访问超出本地或共享内存区域有效范围的地址时，会发生这种情况。
CUDA_EXCEPTION_6 : “Warp 未对齐地址”	精确	Warp 错误	当 Warp 中的任何线程访问本地或共享内存段中未正确对齐的地址时，会发生这种情况。
CUDA_EXCEPTION_7 : “Warp 无效地址空间”	精确	Warp 错误	当 Warp 中的任何线程执行访问该指令不允许的内存空间的指令时，会发生这种情况。
CUDA_EXCEPTION_8 : “Warp 无效 PC”	精确	Warp 错误	当 Warp 中的任何线程将其 PC 推进到 40 位地址空间之外时，会发生这种情况。
CUDA_EXCEPTION_9 : “Warp 硬件堆栈溢出”	精确	Warp 错误	当 Warp 中的任何线程触发硬件堆栈溢出时，会发生这种情况。这种情况应该很少发生。
CUDA_EXCEPTION_10 : “设备非法地址”	精确	全局错误	当线程访问非法（越界）全局地址时，会发生这种情况。
CUDA_EXCEPTION_11 : “已弃用”	已弃用	已弃用	此异常已弃用，应视为 CUDA_EXCEPTION_0。
CUDA_EXCEPTION_12 : “Warp 断言”	精确	每个 Warp	当 Warp 中的任何线程命中设备端断言时，会发生这种情况。
CUDA_EXCEPTION_13 : “已弃用”	已弃用	已弃用	此异常已弃用，应视为 CUDA_EXCEPTION_0。
CUDA_EXCEPTION_14 : “Warp 非法地址”	精确	每个 Warp	当线程访问非法（越界）全局/本地/共享地址时，会发生这种情况。
CUDA_EXCEPTION_15 : “无效托管内存访问”	精确	每个主机线程	当主机线程尝试访问当前由 GPU 使用的托管内存时，会发生这种情况。
CUDA_EXCEPTION_13 : “已弃用”	已弃用	已弃用	此异常已弃用，应视为 CUDA_EXCEPTION_0。
CUDA_EXCEPTION_17 : “集群目标 Block 不存在”	不精确	每个 Cuda 集群	当 Block 中的任何线程访问另一个 Block，而该 Block 超出属于集群的 Block 的有效范围时，会发生这种情况。
CUDA_EXCEPTION_18 : “集群超出范围地址”	不精确	每个 Cuda 集群	当 Block 中的任何线程访问超出属于集群的共享内存区域的有效范围的地址时，会发生这种情况。

10.3. 自动步进

自动步进是一个命令，用于提高 CUDA 异常的精度到精确的 Lane 和指令，否则它们将不会如此精确。

在正常执行下，异常可能会在异常发生后几条指令才报告，或者除非异常是 Lane 错误，否则可能不知道异常发生的精确线程。但是，如果在异常发生时程序正在单步执行，则可以确定异常的精确来源。手动单步执行是一个缓慢而繁琐的过程；单步执行比正常执行花费的时间长得多，并且用户必须单独单步执行每个 Warp。

自动步进通过允许用户指定他们怀疑可能发生异常的代码段来帮助用户，并且在程序运行时，这些代码段会自动且透明地单步执行。程序的其余部分正常执行，以最大程度地减少单步执行引起的减速。如果异常发生在这些代码段中，则将报告异常的精确来源。因此，通过使用自动步进，可以快速且轻松地找到异常发生的精确指令和线程。

注意

autostep 命令是使用断点实现的。如果在当前 PC 上设置了自动步进，则只有在下次遇到 PC 时才会评估自动步进。

自动步进用法

autostep [LOCATION]
autostep [LOCATION] for LENGTH [lines|instructions]

• LOCATION 可以是您用于指定断点位置的任何内容，例如行号、函数名称或以星号开头的指令地址。如果未指定 LOCATION，则使用当前指令地址。

• LENGTH 以行数或指令数指定自动步进窗口的大小（lines 和 instructions 可以缩写，例如 l 或 i）。如果未指定长度类型，则默认为 lines。如果省略 for 子句，则默认为 1 行。

• astep 可以用作 autostep 命令的别名。

• 在自动步进期间对函数的调用将被跳过。

• 在发散的情况下，自动步进窗口的长度由每个 Warp 中第一个活动 Lane 执行的行数或指令数决定。
  发散的 Lane 也被单步执行，但它们执行的指令不计入自动步进窗口的长度。

• 如果在自动步进窗口内发生断点，则命中断点的 Warp 在程序恢复时将不会继续自动步进。但是，其他 Warp 可能会继续自动步进。

• 不支持重叠的自动步进。

如果在执行另一个自动步进时遇到自动步进，则第二个自动步进将被忽略。

如果在内存错误的位置之前设置了自动步进，并且没有命中内存错误，则可能是选择的窗口太小。这可能是由于在自动步进位置的地址和触发内存错误的指令之间存在函数调用而引起的。在这种情况下，要么增加窗口的大小以确保包含错误的指令，要么将自动步进位置移动到更接近错误指令执行时间的指令。

相关命令

自动步进和断点共享相同的编号，因此大多数适用于断点的命令也适用于自动步进。

info autosteps 显示所有断点和自动步进。它类似于 info breakpoints。

(cuda-gdb) info autosteps
Num  Type      Disp Enb Address            What
1    autostep  keep y   0x0000000000401234 in merge at sort.cu:30 for 49 instructions
3    autostep  keep y   0x0000000000489913 in bubble at sort.cu:94 for 11 lines

disable autosteps 禁用自动步进。它等效于 disable breakpoints n。

delete autosteps n 删除自动步进。它等效于 delete breakpoints n。

ignore n i 告诉调试器在调试器进入自动步进 *n* 的窗口 *i* 次后，不要单步执行。此命令已存在于断点中。

11. 演练示例

本章包含三个 CUDA-GDB 演练示例

• 示例：bitreverse

• 示例：autostep

• 示例：MPI CUDA 应用程序

11.1. 示例：bitreverse

本节通过调试一个名为 bitreverse 的示例应用程序来演示 CUDA-GDB 的演练，该应用程序对数据集执行简单的 8 位反转。

源代码

1  #include <stdio.h>
2  #include <stdlib.h>
3
4  // Simple 8-bit bit reversal Compute test
5
6  #define N 256
7
8  __global__ void bitreverse(void *data) {
9     unsigned int *idata = (unsigned int*)data;
10    extern __shared__ int array[];
11
12    array[threadIdx.x] = idata[threadIdx.x];
13
14    array[threadIdx.x] = ((0xf0f0f0f0 & array[threadIdx.x]) >> 4) |
15                        ((0x0f0f0f0f & array[threadIdx.x]) << 4);
16    array[threadIdx.x] = ((0xcccccccc & array[threadIdx.x]) >> 2) |
17                        ((0x33333333 & array[threadIdx.x]) << 2);
18    array[threadIdx.x] = ((0xaaaaaaaa & array[threadIdx.x]) >> 1) |
19                         ((0x55555555 & array[threadIdx.x]) << 1);
20
21    idata[threadIdx.x] = array[threadIdx.x];
22 }
23
24 int main(void) {
25     void *d = NULL; int i;
26     unsigned int idata[N], odata[N];
27
28     for (i = 0; i < N; i++)
29         idata[i] = (unsigned int)i;
30
31     cudaMalloc((void**)&d, sizeof(int)*N);
32     cudaMemcpy(d, idata, sizeof(int)*N,
33                cudaMemcpyHostToDevice);
34
35     bitreverse<<<1, N, N*sizeof(int)>>>(d);
36
37     cudaMemcpy(odata, d, sizeof(int)*N,
38                cudaMemcpyDeviceToHost);
39
40     for (i = 0; i < N; i++)
41        printf("%u -> %u\n", idata[i], odata[i]);
42
43     cudaFree((void*)d);
44     return 0;
45 }

11.1.1. 代码演练

1. 首先，通过在 shell 提示符下输入以下命令来编译用于调试的 bitreverse.cu CUDA 应用程序

   $ nvcc -g -G bitreverse.cu -o bitreverse
   

   此命令假定源文件名为 bitreverse.cu，并且编译不需要额外的编译器标志。另请参见 调试编译

2. 通过在 shell 提示符下输入以下命令来启动 CUDA 调试器

   $ cuda-gdb bitreverse
   

3. 设置断点。在此处设置主机 (main) 和 GPU (bitreverse) 断点。此外，在设备函数中的特定行 (bitreverse.cu:18) 设置断点。

   (cuda-gdb) break main
   Breakpoint 1 at 0x18e1: file bitreverse.cu, line 25.
   (cuda-gdb) break bitreverse
   Breakpoint 2 at 0x18a1: file bitreverse.cu, line 8.
   (cuda-gdb) break 21
   Breakpoint 3 at 0x18ac: file bitreverse.cu, line 21.
   

4. 运行 CUDA 应用程序，它将执行直到到达上一步中设置的第一个断点 (main)。

   (cuda-gdb) run
   Starting program: /Users/CUDA_User1/docs/bitreverse
   Reading symbols for shared libraries
   ..++........................................................... done
   
   Breakpoint 1, main () at bitreverse.cu:25
   25  void *d = NULL; int i;
   

5. 此时，可以输入命令以推进执行或打印程序状态。对于此演练，让我们继续直到设备 Kernel 启动。

   (cuda-gdb) continue
   Continuing.
   Reading symbols for shared libraries .. done
   Reading symbols for shared libraries .. done
   [Context Create of context 0x80f200 on Device 0]
   [Launch of CUDA Kernel 0 (bitreverse<<<(1,1,1),(256,1,1)>>>) on Device 0]
   Breakpoint 3 at 0x8667b8: file bitreverse.cu, line 21.
   [Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
   
   Breakpoint 2, bitreverse<<<(1,1,1),(256,1,1)>>> (data=0x110000) at bitreverse.cu:9
   9   unsigned int *idata = (unsigned int*)data;
   

   CUDA−GDB 已检测到已到达 CUDA 设备 Kernel。调试器打印当前焦点的 CUDA 线程。

6. 使用 info cuda threads 命令验证焦点的 CUDA 线程，并在主机线程和 CUDA 线程之间切换

   (cuda-gdb) info cuda threads
     BlockIdx ThreadIdx To BlockIdx ThreadIdx Count            Virtual PC
   Filename   Line
   
   Kernel 0
   *  (0,0,0)    (0,0,0)    (0,0,0) (255,0,0)    256 0x0000000000866400 bitreverse.cu     9
   (cuda-gdb) thread
   [Current thread is 1 (process 16738)]
   (cuda-gdb) thread 1
   [Switching to thread 1 (process 16738)]
   #0  0x000019d5 in main () at bitreverse.cu:34
   34    bitreverse<<<1, N, N*sizeof(int)>>>(d);
   (cuda-gdb) backtrace
   #0  0x000019d5 in main () at bitreverse.cu:34
   (cuda-gdb) info cuda kernels
   Kernel Dev Grid   SMs Mask GridDim  BlockDim        Name Args
        0   0    1 0x00000001 (1,1,1) (256,1,1) bitreverse data=0x110000
   (cuda-gdb) cuda kernel 0
   [Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
   9    unsigned int *idata = (unsigned int*)data;
   (cuda-gdb) backtrace
   #0   bitreverse<<<(1,1,1),(256,1,1)>>> (data=0x110000) at bitreverse.cu:9
   

7. 通过打印 Block 和线程索引来证实此信息

   (cuda-gdb) print blockIdx
   $1 = {x = 0, y = 0}
   (cuda-gdb) print threadIdx
   $2 = {x = 0, y = 0, z = 0)
   

8. 还可以打印 Grid 和 Block 维度

   (cuda-gdb) print gridDim
   $3 = {x = 1, y = 1}
   (cuda-gdb) print blockDim
   $4 = {x = 256, y = 1, z = 1)
   

9. 推进 Kernel 执行并验证一些数据

   (cuda-gdb) next
   12       array[threadIdx.x] = idata[threadIdx.x];
   (cuda-gdb) next
   14       array[threadIdx.x] = ((0xf0f0f0f0 & array[threadIdx.x]) >> 4) |
   (cuda-gdb) next
   16       array[threadIdx.x] = ((0xcccccccc & array[threadIdx.x]) >> 2) |
   (cuda-gdb) next
   18       array[threadIdx.x] = ((0xaaaaaaaa & array[threadIdx.x]) >> 1) |
   (cuda-gdb) next
   
   Breakpoint 3, bitreverse <<<(1,1),(256,1,1)>>> (data=0x100000) at bitreverse.cu:21
   21             idata[threadIdx.x] = array[threadIdx.x];
   (cuda-gdb) print array[0]@12
   $7 = {0, 128, 64, 192, 32, 160, 96, 224, 16, 144, 80, 208}
   (cuda-gdb) print/x array[0]@12
   $8 = {0x0, 0x80, 0x40, 0xc0, 0x20, 0xa0, 0x60, 0xe0, 0x10, 0x90, 0x50,
   0xd0}
   
   (cuda-gdb) print &data
   $9 = (@global void * @parameter *) 0x10
   (cuda-gdb) print *(@global void * @parameter *) 0x10
   $10 = (@global void * @parameter) 0x100000
   

   结果输出取决于内存位置的当前内容。

10. 由于线程 (0,0,0) 反转 0 的值，因此切换到不同的线程以显示更有趣的数据

    (cuda-gdb) cuda thread 170
    [Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread
    (170,0,0), device 0, sm 0, warp 5, lane 10]
    

11. 删除断点并继续程序直到完成

    (cuda-gdb) delete breakpoints
    Delete all breakpoints? (y or n) y
    (cuda-gdb) continue
    Continuing.
    
    Program exited normally.
    (cuda-gdb)
    

11.2. 示例：autostep

本节介绍如何使用自动步进命令，并演示它如何帮助提高内存错误报告的精度。

源代码

1  #define NUM_BLOCKS 8
2  #define THREADS_PER_BLOCK 64
3
4  __global__ void example(int **data) {
5    int value1, value2, value3, value4, value5;
6    int idx1, idx2, idx3;
7
8    idx1 = blockIdx.x * blockDim.x;
9    idx2 = threadIdx.x;
10   idx3 = idx1 + idx2;
11   value1 = *(data[idx1]);
12   value2 = *(data[idx2]);
13   value3 = value1 + value2;
14   value4 = value1 * value2;
15   value5 = value3 + value4;
16   *(data[idx3]) = value5;
17   *(data[idx1]) = value3;
18   *(data[idx2]) = value4;
19   idx1 = idx2 = idx3 = 0;
20 }
21
22 int main(int argc, char *argv[]) {
23   int *host_data[NUM_BLOCKS * THREADS_PER_BLOCK];
24   int **dev_data;
25   const int zero = 0;
26
27   /* Allocate an integer for each thread in each block */
28   for (int block = 0; block < NUM_BLOCKS; block++) {
29     for (int thread = 0; thread < THREADS_PER_BLOCK; thread++) {
30       int idx = thread + block * THREADS_PER_BLOCK;
31       cudaMalloc(&host_data[idx], sizeof(int));
32       cudaMemcpy(host_data[idx], &zero, sizeof(int),
33                  cudaMemcpyHostToDevice);
34     }
35   }
36
37   /* This inserts an error into block 3, thread 39*/
38   host_data[3*THREADS_PER_BLOCK  + 39] = NULL;
39
40   /* Copy the array of pointers to the device */
41   cudaMalloc((void**)&dev_data,  sizeof(host_data));
42   cudaMemcpy(dev_data, host_data, sizeof(host_data), cudaMemcpyHostToDevice);
43
44   /* Execute example */
45   example <<< NUM_BLOCKS, THREADS_PER_BLOCK >>> (dev_data);
46   cudaThreadSynchronize();
47 }

在这个小示例中，我们有一个整数指针数组，我们想对整数进行一些操作。但是，假设其中一个指针为 NULL，如第 38 行所示。当我们尝试访问与 Block 3、线程 39 对应的整数时，这将导致抛出 CUDA_EXCEPTION_10 "Device Illegal Address" 异常。当我们尝试写入该值时，此异常应发生在第 16 行。

11.2.1. 使用自动步进进行调试

1. 像往常一样编译示例并启动 CUDA−GDB。我们首先运行程序

   (cuda-gdb) run
   Starting program: /home/jitud/cudagdb_test/autostep_ex/example
   [Thread debugging using libthread_db enabled] [New Thread 0x7ffff5688700 (LWP 9083)]
   [Context Create of context 0x617270 on Device 0]
   [Launch of CUDA Kernel 0 (example<<<(8,1,1),(64,1,1)>>>) on Device 0]
   
   Program received signal CUDA_EXCEPTION_10, Device Illegal Address.
   [Switching focus to CUDA kernel 0, grid 1, block (1,0,0), thread (0,0,0), device 0, sm 1, warp 0, lane 0]
   0x0000000000796f60 in example (data=0x200300000) at example.cu:17
   17        *(data[idx1]) = value3;
   

   正如预期的那样，我们收到了 CUDA_EXCEPTION_10。但是，报告的线程是 Block 1，线程 0，行号是 17。由于 CUDA_EXCEPTION_10 是全局错误，因此没有报告线程信息，因此我们将不得不手动检查所有 512 个线程。

2. 设置 autosteps。为了获得更准确的信息，我们推断由于 CUDA_EXCEPTION_10 是内存访问错误，因此它必须发生在访问内存的代码上。这发生在第 11、12、16、17 和 18 行，因此我们为这些区域设置了两个自动步进窗口

   (cuda-gdb) autostep 11 for 2 lines
   Breakpoint 1 at 0x796d18: file example.cu, line 11.
   Created autostep of length 2 lines
   (cuda-gdb) autostep 16 for 3 lines
   Breakpoint 2 at 0x796e90: file example.cu, line 16.
   Created autostep of length 3 lines
   

3. 最后，我们使用这些自动步进再次运行程序

   (cuda-gdb) run
   The program being debugged has been started already.
   Start it from the beginning? (y or n) y
   [Termination of CUDA Kernel 0 (example<<<(8,1,1),(64,1,1)>>>) on Device 0]
   Starting program: /home/jitud/cudagdb_test/autostep_ex/example
   [Thread debugging using libthread_db enabled]
   [New Thread 0x7ffff5688700 (LWP 9089)]
   [Context Create of context 0x617270 on Device 0]
   [Launch of CUDA Kernel 1 (example<<<(8,1,1),(64,1,1)>>>) on Device 0]
   [Switching focus to CUDA kernel 1, grid 1, block (0,0,0), thread (0,0,0),
   device 0, sm 0, warp 0, lane 0]
   
   Program received signal CUDA_EXCEPTION_10, Device Illegal Address.
   [Current focus set to CUDA kernel 1, grid 1, block (3,0,0), thread
   (32,0,0), device 0, sm 1, warp 3, lane 0]
   Autostep precisely caught exception at example.cu:16 (0x796e90)
   

   这次我们正确地在第 16 行捕获了异常。即使 CUDA_EXCEPTION_10 是全局错误，我们也已将其范围缩小到 Warp 错误，因此我们现在知道抛出异常的线程必须与 Block 3、线程 32 处于同一 Warp 中。

在此示例中，我们仅通过设置两个 autosteps 并重新运行程序，就将错误的范围从 512 个线程缩小到 32 个线程。

11.3. 示例：MPI CUDA 应用程序

对于大规模 MPI CUDA 应用程序调试，NVIDIA 建议使用我们的合作伙伴 Allinea 和 Totalview 提供的并行调试器。两者都是优秀的并行调试器，具有对 CUDA 的扩展支持。但是，对于调试较小的应用程序，或者仅调试大型应用程序中的少数进程，可以使用 CUDA-GDB。

如果集群节点具有 xterm 支持，请以与使用作业启动器启动 gdb 相同的方式启动 CUDA-GDB。例如

$ mpirun -np 4 -host nv1,nv2 xterm -e cuda-gdb a.out

您可能必须导出 DISPLAY 变量以确保 xterm 找到返回到您的显示器的路径。例如

$ mpirun -np 4 -host nv1,nv2 -x DISPLAY=host.nvidia.com:0 xterm -e cuda-gdb a.out

作业启动器有不同的方式将环境变量导出到集群节点。有关更多详细信息，请查阅您的作业启动器文档。

如果您的集群环境不支持 xterm，您可以程序内部插入一个自旋锁循环，ssh 登录到计算节点，并附加到 MPI 进程。在程序开始附近的某个位置，添加类似于以下代码片段的代码：

{
    int i = 0;
    char host[256];
    printf("PID %d on node %s is ready for attach\n",
            getpid(), host);
    fflush(stdout);
    while (0 == i) {
        sleep(5);
    }
}

重新编译并启动应用程序。启动后，ssh 登录到感兴趣的节点，并使用 CUDA-GDB 附加到进程。将变量 i 设置为 1 以跳出循环

$ mpirun -np 2 -host nv1,nv2 a.out
PID 20060 on node nv1 is ready for attach
PID 5488 on node nv2 is ready for attach

$ ssh nv1
[nv1]$ cuda-gdb --pid 5488

$ ssh nv2
[nv2]$ cuda-gdb --pid 20060

对于更大的应用程序，您可以根据 MPI 进程的 rank 使用 MPI_Comm_rank 函数来条件化自旋锁循环。

对于计算能力低于 6.0 的设备，多个调试器 中描述的软件抢占解决方法不适用于 MPI 应用程序。对于这些 GPU，请确保每个 MPI 进程都以唯一的 GPU 为目标。

如果设置了 CUDA_VISIBLE_DEVICES，则可能会导致 MPI 应用程序中的 GPU 选择逻辑出现问题。它也可能阻止节点上 GPU 之间的 CUDA IPC 工作。

12. 提示与技巧

本节作为 CUDA-GDB 用户可以利用的，但在其他地方未记录的高级设置以及各种提示和技巧的参考。

12.1. set cuda break_on_launch

要在每个启动的内核的第一条指令处中断，请将 break_on_launch 选项设置为 application

(cuda-gdb) set cuda break_on_launch application

可能的选项包括

none

无 Kernel，应用程序或系统（默认）

application

用户应用程序启动的 Kernel

system

驱动程序启动的任何 Kernel，例如 memset

all

任何 Kernel、应用程序和系统

这些自动断点不会通过 info breakpoints 命令显示，并且与单个断点分开管理。关闭此选项不会删除设置为相同地址的其他单个断点，反之亦然。

12.2. set cuda launch_blocking

启用后，内核启动是同步的，就像已将环境变量 CUDA_LAUNCH_BLOCKING 设置为 1 一样。一旦阻塞，启动实际上将被序列化，并且可能更容易调试。

• (cuda-gdb) set cuda launch_blocking off
  

  内核启动是同步或异步的，由应用程序决定。这是默认设置。

• (cuda-gdb) set cuda launch_blocking on
  

  内核启动是同步的。如果应用程序已经启动，则更改仅在当前会话终止后生效。

12.3. set cuda notify

每当发生 CUDA 事件时，都需要通知调试器。通知以信号的形式发送到主机线程。接收该特殊信号的主机线程由 set cuda notify 选项确定。

• (cuda-gdb) set cuda notify youngest
  

  线程 ID 最小的主机线程将接收通知信号（默认）。

• (cuda-gdb) set cuda notify random
  

  任意主机线程将接收通知信号。

12.4. set cuda ptx_cache

在访问变量的值之前，调试器会检查变量在当前 PC 上是否处于活动状态。在 CUDA 设备上，变量可能并非始终处于活动状态，并且将被报告为“已优化掉”。

CUDA-GDB 提供了一个选项，通过在 PTX 寄存器级别缓存变量的值来规避此限制。每个源变量都编译成一个 PTX 寄存器，然后映射到一个或多个硬件寄存器。使用编译器发出的调试信息，调试器可能能够根据上次映射到的硬件寄存器缓存 PTX 寄存器的值。

此优化始终是正确的。启用后，缓存的值将显示为从实际硬件寄存器读取的正常值，并带有 (cached) 前缀。该优化仅在单步执行代码时生效。

• (cuda-gdb) set cuda ptx_cache off
  

  调试器仅读取活动变量的值。

• (cuda-gdb) set cuda ptx_cache on
  

  调试器将在可能的情况下使用缓存的值。此设置是默认设置，并且始终是安全的。

12.5. set cuda single_stepping_optimizations

单步执行可能需要花费大量时间。启用后，此选项会告知调试器使用安全技巧来加速单步执行。

• (cuda-gdb) set cuda single_stepping_optimizations off
  

  调试器不会尝试加速单步执行。这是 5.5 版本及更早版本中的唯一且默认行为。

• (cuda-gdb) set cuda single_stepping_optimizations on
  

  调试器将使用安全技术来加速单步执行。这是从 6.0 版本开始的默认设置。

12.6. set cuda thread_selection

当调试器必须选择一个活动线程来关注时，该决定由启发式方法指导。set cuda thread_selection 指导这些启发式方法。

• (cuda-gdb) set cuda thread_selection logical
  

  选择具有最低 blockIdx/threadIdx 坐标的线程。

• (cuda-gdb) set cuda thread_selection physical
  

  选择具有最低 dev/sm/warp/lane 坐标的线程。

12.7. set cuda value_extrapolation

在访问变量的值之前，调试器会检查变量在当前 PC 上是否处于活动状态。在 CUDA 设备上，变量可能并非始终处于活动状态，并且将被报告为“已优化掉”。

CUDA-GDB 提供了一个选项，通过在调试器原本会将其标记为已优化掉时，机会性地外推变量的值来规避此限制。外推不能保证准确，必须谨慎使用。如果用于存储变量值的寄存器自上次变量被视为活动状态以来已被重用，则报告的值将是错误的。因此，使用该选项打印的任何值都将标记为 "(possibly)"。

• (cuda-gdb) set cuda value_extrapolation off
  

  调试器仅读取活动变量的值。此设置是默认设置，并且始终是安全的。

• (cuda-gdb) set cuda value_extrapolation on
  

  调试器将尝试外推变量的值，使其超出各自的活动范围。此设置可能会报告错误的值。

12.8. 调试 Docker 容器

在 Docker 容器中调试应用程序时，需要启用 PTRACE 功能。用户还需要确保根文件系统同时设置了读/写权限。

要启用 PTRACE 功能，请将以下内容添加到您的 Docker run 命令中

--cap-add=SYS_PTRACE

12.9. 切换到经典调试器后端

在 Linux 平台上，CTK 11.8 版本引入了一个名为统一调试器 (UD) 的新调试器后端。UD 允许与诸如 cuda-gdb 和 NVIDIA® Nsight™ VSE 等调试工具共享统一的调试器后端。UD 跨多个平台（包括 Windows 和 Linux）受支持。最终用户使用 UD 的体验对于现有工具的使用是透明的。

之前的调试器后端，称为经典调试器后端，仍然可以通过在启动 CUDA-GDB 之前在环境中将 CUDBG_USE_LEGACY_DEBUGGER 设置为 1 来使用。

Maxwell GPU 不支持 UD。用户必须切换到经典调试器后端才能在 Maxwell GPU 上调试其应用程序。

12.10. 线程块集群

使用线程块集群的 CUDA 应用程序将在 CUDA 焦点中看到显示的集群索引。集群索引和集群维度都可以通过打印便捷变量 clusterIdx 和 clusterDim 来查询。

12.11. 调试 OptiX/RTCore 应用程序

调试使用 OptiX/RTCore 构建的程序时，可能需要将环境变量 OPTIX_FORCE_DEPRECATED_LAUNCHER 设置为 1。如果断点无法命中，请尝试在启动应用程序之前设置此环境变量。

12.12. 在 Windows Subsystem for Linux 上调试

如果您无法在 Windows Subsystem for Linux 上使用调试器，请确保通过将注册表项 >HKEY_LOCAL_MACHINE\SOFTWARE\NVIDIA Corporation\GPUDebugger\EnableInterface 设置为 (DWORD) 1 来启用调试接口

12.13. 从主机线程访问非托管设备内存

在访问通过 CUDA 内存分配 API 创建的非托管设备内存时，需要特别考虑以下事项

CUdeviceptr d_global;
cuMemAlloc(&d_global, sizeof(uint32_t));

uint32_t h_global_32 = 0x1234;
cuMemcpyHtoD(d_global, &h_global, sizeof(uint32_t));

当专注于主机线程时直接检查 d_global 的内容时，我们将观察到不正确的值

(cuda-gdb) print /x *d_global
$1 = 0x0

用户需要显式类型转换为正确的地址空间标识符，以检查设备上的非托管内存

(cuda-gdb) print *(@global unsigned long long *)d_global
$2 = 0x1234

当专注于 CUDA 线程时，显式转换是不必要的。

13. 支持的平台

主机平台要求

CUDA-GDB 在与其一起提供的 CUDA 工具包支持的所有平台上均受支持。有关更多信息，请参阅 CUDA 工具包发行说明。

GPU 要求

当前 CUDA 版本支持的所有 CUDA 功能 GPU 都支持调试。

GDB Python 集成

cuda-gdb 中的 GDB Python 集成通过多构建机制支持，以便跨不同平台支持多个 python3 解释器。cuda-gdb 程序是一个 shell 脚本，它根据系统上可用的 python 版本选择关联的支持的 cuda-gdb 二进制文件。支持以下 Python 版本：Python 3.8、Python 3.9、Python 3.10、Python 3.11 和 Python 3.12

Windows Subsystem for Linux (WSL)

• cuda-gdb 支持在 WSL2 上调试 CUDA 应用程序。

• 确保通过将注册表项 >HKEY_LOCAL_MACHINE\SOFTWARE\NVIDIA Corporation\GPUDebugger\EnableInterface 设置为 (DWORD) 1 来启用此功能。

• 调试计算密集型应用程序可能需要增加或禁用 TDR。

14. 支持的操作系统上的常见问题

以下是当前版本在支持的操作系统上的已知问题以及如何修复它们。

Python 未初始化

发生这种情况的原因是机器上缺少 Python 3.x 库，安装它可以解决此问题。这也可能是由于安装的 libpython 的 major.minor 版本与 PATH 中的默认 python3 解释器不匹配造成的。必须提供与 PATH 中的默认 python3 解释器匹配的 libpython 版本。可以使用 python3 --version 命令确定 libpython 版本。例如，以下命令将告诉我们需要在默认库搜索路径中安装 libpython3.8.so*

$ python3 --version
Python 3.8.10

以下是安装正确 libpython 的具体命令。

RHEL 8/9

$ sudo yum -y install python3-libs

Debian 10/11/12

$ sudo apt-get -y install libpython3-stdlib

Fedora 39

$ sudo yum -y install python3-libs

OpenSUSE 15

$ sudo zypper install -y libpython3

Ubuntu 20.04/22.04

$ sudo apt-get -y install python3.8 $ sudo apt-get -y install libpython3.8

15. 已知问题

以下是当前版本的已知问题。

• Windows 或 WSL 上不支持启用 MCDM 的 Hopper 架构上的调试。

• 在加载 __device__ 或 __global__ 函数的模块之前，在其内部的行上设置断点可能会导致断点临时设置在源代码中下面函数的首行。一旦加载了目标函数的模块，断点将正确重置。同时，断点可能会被命中，具体取决于应用程序。在这些情况下，可以安全地忽略断点，并且可以恢复应用程序。

• scheduler-locking 选项不能设置为 on。

• 从内核单步跳出后再次单步执行会导致不确定的行为。建议改用“continue”命令。

• 不支持在 cuda-gdb 中附加到启用了软件抢占的 CUDA 应用程序。

• 不支持在 QNX 上附加到 CUDA 应用程序。

• 不支持在 MPS 客户端模式下附加到 CUDA 应用程序。

• 不支持使用 cuda-gdb 附加到 MPS 服务器进程 (nvidia-cuda-mps-server)，或使用 cuda-gdb 启动 MPS 服务器。

• 如果 CUDA 应用程序在 MPS 客户端模式下使用 cuda-gdb 启动，则 MPS 客户端将等待直到所有其他 MPS 客户端终止，然后将作为非 MPS 应用程序运行。

• 当调试器单步跳过内联例程时，会发生明显的性能下降。

由于内联代码块可能具有多个退出点，因此在后台，调试器会单步执行每个指令，直到到达退出点，这对于大型例程会产生相当大的成本。建议采取以下操作以避免此问题

• 声明函数时避免使用 __forceinline__。（对于使用调试信息编译的代码，只有使用 __forceinline__ 关键字声明的例程实际上是内联的）

• 使用 until <行号> 命令单步跳过内联子例程。

• 在 Jetson 上，调用 cuda API 可能会导致调试器跳转到 _dl_catch_exception()。解决方法是继续。

• 在 Jetson 和 Drive 设备上，只有在以 root 权限运行调试器时，GPU 调试才能正常工作。需要更改 devfs 节点权限才能在不以 root 身份运行的情况下使调试器工作。

• 如果诱导陷阱 (__trap()) 是设备从断点恢复后执行的下一条指令，则调试器可能会错过报告。

• 如果先前为空的 SM 上启动了新的 warp，则调试器可能会错过在恢复期间报告断点或异常。

• 调试器使用系统上安装的 libpython。使用 Python 脚本功能会将 cuda-gdb 暴露于与系统 libpython 版本中相同的漏洞。建议始终保持系统 libpython 库为最新版本。

• 调试器不支持访问使用 CUDA IPC API 从其他进程导入的共享内存分配。调试器尝试访问这些共享内存分配将导致错误，指出不支持访问通过 IPC 共享的内存分配。

• break_on_launch 将无法与 OptiX/RTCore 程序一起使用，除非将 OPTIX_FORCE_DEPRECATED_LAUNCHER 设置为 1。

• 在 QNX 上，需要将 QNX_TARGET 环境变量设置为指向目标 rootfs。如果此变量不可用，请将环境变量设置为空字符串。

16. 声明

16.1. 声明

本文档仅供参考，不得视为对产品的特定功能、条件或质量的保证。NVIDIA Corporation（“NVIDIA”）对本文档中包含的信息的准确性或完整性不作任何明示或暗示的陈述或保证，并且对本文档中包含的任何错误不承担任何责任。NVIDIA 对因使用此类信息或因使用此类信息而可能导致的任何专利或第三方其他权利的侵权行为的后果或使用不承担任何责任。本文档不构成对开发、发布或交付任何材料（如下定义）、代码或功能的承诺。

NVIDIA 保留随时更正、修改、增强、改进和对本文档进行任何其他更改的权利，恕不另行通知。

客户应在下订单前获取最新的相关信息，并应验证此类信息是否为最新且完整。

除非 NVIDIA 和客户的授权代表签署的个别销售协议（“销售条款”）另有约定，否则 NVIDIA 产品的销售受订单确认时提供的 NVIDIA 标准销售条款和条件的约束。NVIDIA 在此明确反对将任何客户一般条款和条件应用于购买本文档中引用的 NVIDIA 产品。本文档未直接或间接形成任何合同义务。

NVIDIA 产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命支持设备，也不适用于 NVIDIA 产品的故障或故障可以合理预期会导致人身伤害、死亡或财产或环境损害的应用。NVIDIA 对 NVIDIA 产品在此类设备或应用中的包含和/或使用不承担任何责任，因此，此类包含和/或使用由客户自行承担风险。

NVIDIA 不保证或声明基于本文档的产品将适用于任何特定用途。NVIDIA 不一定对每种产品的所有参数进行测试。客户全权负责评估和确定本文档中包含的任何信息的适用性，确保产品适合并符合客户计划的应用，并为应用执行必要的测试，以避免应用或产品的默认设置。客户产品设计中的缺陷可能会影响 NVIDIA 产品的质量和可靠性，并可能导致超出本文档中包含的其他或不同的条件和/或要求。对于可能基于以下原因或归因于以下原因的任何默认设置、损坏、成本或问题，NVIDIA 不承担任何责任：(i) 以任何违反本文档的方式使用 NVIDIA 产品，或 (ii) 客户产品设计。

本文档未授予 NVIDIA 专利权、版权或本文档项下的其他 NVIDIA 知识产权的任何明示或暗示的许可。NVIDIA 发布的有关第三方产品或服务的信息不构成 NVIDIA 授予使用此类产品或服务的许可，也不构成对其的保证或认可。使用此类信息可能需要获得第三方的专利或第三方其他知识产权的许可，或获得 NVIDIA 的专利或 NVIDIA 其他知识产权的许可。

仅当事先获得 NVIDIA 的书面批准，未经修改且完全符合所有适用的出口法律和法规进行复制，并附带所有相关的条件、限制和声明时，才允许复制本文档中的信息。

本文档和所有 NVIDIA 设计规范、参考板、文件、图纸、诊断程序、列表和其他文档（统称为“材料”）均按“原样”提供。NVIDIA 对材料不作任何明示、暗示、法定或其他形式的保证，并明确否认所有关于不侵权、适销性和特定用途适用性的暗示保证。在法律未禁止的范围内，在任何情况下，NVIDIA 均不对因使用本文档而引起的任何损害（包括但不限于任何直接、间接、特殊、偶然、惩罚性或后果性损害，无论因何种原因引起，也无论责任理论如何）承担责任，即使 NVIDIA 已被告知可能发生此类损害。尽管客户可能因任何原因而遭受任何损害，但 NVIDIA 对本文所述产品的客户的累计责任应根据产品的销售条款进行限制。

16.2. OpenCL

OpenCL 是 Apple Inc. 的商标，已获得 Khronos Group Inc. 的许可使用。

16.3. 商标

NVIDIA 和 NVIDIA 徽标是 NVIDIA Corporation 在美国和其他国家/地区的商标或注册商标。其他公司和产品名称可能是与其关联的各自公司的商标。