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是GNU项目调试器GDB的扩展。该工具为开发人员提供了一种调试在实际硬件上运行的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）编译的内核，同时也支持调试运行时编译的内核，这被称为即时编译（JIT compilation）。

1.3. 关于本文档

本文档是CUDA-GDB的主要文档，其结构更类似于用户手册而非参考手册。文档后续部分将介绍如何安装和使用CUDA-GDB来调试CUDA内核，以及如何使用新增到GDB中的CUDA命令。同时提供了一些操作示例。假定用户已掌握用于调试主机应用程序的基本GDB命令。

2. 版本说明

12.8 版本发布

Maxwell, Pascal, and Volta deprecation notice

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

Coredump env var deprecation notice

• 在CUDA 12.5中，环境变量CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION和CUDA_ENABLE_LIGHTWEIGHT_COREDUMP已被弃用，改用CUDA_COREDUMP_GENERATION_FLAGS。

Features

• 新增解析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 类型。

Fixed Issues

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

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

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

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

12.7 版本发布

Features

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

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

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

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

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

• 添加支持，用于在遇到集群异常时识别目标块。

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

Fixed Issues

• 修复了因遇到无效的CUDA寄存器编号导致CUDA寄存器外推失败而引发的崩溃问题。

• 修复了在存在活跃CUDA次级进程时无法切换到cudacore目标的问题。

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

12.6 版本发布

Features

• 新增 set/show cuda step_divergent_lanes 命令，用于控制当聚焦的CUDA线程不再活跃时对分支线程的自动步进操作。默认开启以保持现有行为。

• 提升了打开包含大量cubin文件的GPU核心文件的性能。

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

• 当开始执行分支线程且step_divergent_lanes开启时发出警告。

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

Fixed Issues

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

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

• 增强了对遇到包含损坏扩展行信息的cubin文件时的错误处理能力。

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

12.5 版本发布

Updated GDB version

• 从GDB 13.1迁移至13.2版本。详见GDB 13.2变更内容

Support removal notice

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

• 已移除对Android的支持。

• 已不再支持 Python 3.6 和 3.7 版本。

Features

• 支持在所有兼容平台上运行原生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命令。

Fixed Issues

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

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

• 修复了同一内核连续启动时遗漏break_on_launch断点的问题。

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

• 修复了在QNX系统上，当cuda-gdbserver数据包乱序到达时导致的崩溃问题。

• 在读取CUDA反汇编代码时遇到错误时提供更好的错误处理。

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

12.4 版本发布

Updated GDB version

• 从GDB 12.1迁移至13.1版本。详见GDB 13.1变更说明

Android deprecation notice

• 对Android的支持已被弃用，将在未来的版本中移除。

Python 3.6 and 3.7 deprecation notice

• 对已终止支持的Python 3.6和3.7版本的支持已被弃用，将在后续版本中移除。

Features

• 性能优化，减少了CUDA调试器API调用的总次数。

• 在使用大量GPU寄存器加载带有设备函数的大型cubin时，性能得到提升。

• 单步跨越warp范围屏障时的性能增强。

• 新增支持从GPU核心转储中打印常量存储区内的值。

Fixed Issues

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

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

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

• 修复了在 inferior teardown 过程中 switch_to_thread_1 的断言问题。

• 修复了在较新型号英特尔处理器上遇到的附加失败问题。

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

12.3 版本发布

macOS host client deprecation notice

• 对macOS主机客户端的CUDA-GDB支持已弃用，将在后续版本中移除。

Features

• 新增支持打印常量存储区中的值。新增便捷函数$_cuda_const_bank(bank, offset)用于获取常量存储区偏移地址。详见Const banks。

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

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

Fixed Issues

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

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

• 新增支持在设备上没有有效线程束时打开GPU核心转储文件。

• 添加了OptiX使用的缺失DWARF操作符。

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

• 修复了在没有CUDA集群时仍显示CUDA集群坐标的问题。

12.2 版本发布

Features

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

• 启用了在开发工具模式下使用机密计算模式进行调试的支持。有关如何启用该模式的更多详细信息，请参阅机密计算部署指南。

Fixed Issues

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

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

• 修复了调试LLVM生成的cubins时名称混淆的问题。

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

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

12.1 版本发布

CUDA Driver API added for controlling core dump behavior

• CTK 12.1和r530驱动程序新增了API，允许开发者在应用程序内部以编程方式启用/配置核心转储设置，而无需使用环境变量。更多信息请参阅CUDA Driver API手册。

Features

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

• 新增支持包含大量节区（超过32767个）的ELF格式cubin文件。

• 为CUDA Graphs添加了break_on_launch支持。

Fixed Issues

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

• 在QNX上修复了一个问题，即info threads错误地报告了已终止的主机线程。

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

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

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

• 修复了在解析自引用结构的DWARF时出现的问题。

12.0 版本发布

Updated GDB version

• 从GDB 10.2迁移至12.1版本。详见GDB 12.1变更说明

Texture and surface reference support removed

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

CUDA Memory Checker integration removed

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

Debugging of applications using CUDA Dynamic Parallelism

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

Features

• 从基础gdb/10.2迁移至gdb/12.1版本。

• 新增了对线程块集群的初步支持。

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

Fixed Issues

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

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

• 修改了内部CUDA动态并行检测断点，仅当启用break_on_launch时才会设置。

• 移除了不支持的gpu_busy_check setting设置。

11.8 版本发布

Features

• 默认使用新的统一调试器(UD)作为调试后端。

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

• 调试器现已在Windows子系统Linux(WSL)上启用。

• 增加了对打印FP8值(E4M3和E5M2)的基本类型支持。

Notes

• 默认情况下，CUDA-GDB将使用新的统一调试器(UD)后端。对于使用Pascal或更新显卡的大多数用户来说，这一变化是无感知的。如需调试Maxwell架构，或强制使用旧版经典调试后端，请在环境变量中将CUDBG_USE_LEGACY_DEBUGGER设置为1。

• 此版本不支持在GH100平台上使用WSL。

11.7 版本发布

Features

• 针对break_on_launch功能进行了重大性能优化，采用新的KERNEL_READY通知机制替代手动设置断点的方式。

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

• 新增了--disable-python命令行选项，用于禁用Python解释器的dlopen功能。

Fixed Issues

• 修复了跟随派生子进程的问题，以避免父进程和子进程同时使用CUDA时出现挂起行为。

• 修复了某些libpython版本中因缺少libpython函数的dlsym而导致的错误。

11.6 版本发布

Updated GDB version

• 从GDB 10.1迁移至10.2版本。详见GDB 10.2变更内容

Features

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

Fixed Issues

• 修复了行信息帧以正确显示源文件名。

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

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

• 修复了cuda-gdbserver初始化检查问题，该问题曾导致QNX无法启动。

11.5 版本发布

Python 3 support on Jetson and Drive Tegra devices

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

Fixed Issues

• 在动态加载libpython3库时增加了健壮的版本检查。加载的libpython3将与PATH中python3运行时的版本匹配。

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

• 新增支持动态加载Python 3.9。

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

11.4 更新 1 版本发布

Known Issues with Fedora 34

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

Fixed Issues

• 为ppc64le和aarch64 SBSA架构启用了Python集成。

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

• 修复了通过cuda-gdbserver进行远程调试时偶发的挂起问题。

• 修复了设置cuda api_failures stop在失败时不触发断点的bug。

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

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

• OpenMP: 增强了对并行区域内打印局部变量的支持。

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

• 修复了在cuda核心调试中选择主机和设备线程焦点的问题。

• 针对QNX远程调试的多项修复。

11.4 版本发布

Updated GDB version

• 从GDB 8.3迁移至10版本（基于GDB 10.1）。查看GDB 10.1变更

Python 3 support

• 不再支持Python 2。CUDA-GDB现在支持Python 3。

GDB TUI mode disabled

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

Kepler deprecation notice

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

Coredump support

• 新增支持通过CUDA_COREDUMP_FILE将核心转储写入命名管道。

Fixed Issues

• 新增支持在核心转储中显示SIGTRAP异常。

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

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

• 修复了创建CUDA临时目录时偶发的竞态条件问题。

• 针对QNX远程调试的多项修复。

11.3 版本发布

Python 2 deprecation notice

• 对Python 2的支持将被弃用。CUDA-GDB将在即将发布的版本中转向构建支持Python 3。

Fixed Issues

• 改进了远程调试的后期附加功能。

11.2 更新 1 版本发布

GDB TUI deprecation notice

• 对GDB TUI模式的支持将被弃用。这将避免在缺乏ncurses-5.5支持的操作系统上出现跨平台依赖不匹配的问题。GDB TUI模式将在即将发布的版本中被禁用。

Fixed Issues

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

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

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

• 修复了启动时因DWARF dies缺少名称导致的段错误。

• 修复了当CUDA内核调用assert时出现的段错误问题。

• 修复了一个导致无法调试大于2GB的cubin文件的错误。

• 为使用--lineinfo编译的cubins添加了小幅可用性改进。

• 修复了在CLion中使用CUDA-GDB时由格式化打印导致的段错误问题。

11.1 版本发布

Updated GDB version

• 从GDB 8.2迁移至8.3版本（基于gdb 8.3.1）。查看gdb 8.3.1变更内容

Support for SM 8.6

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

Updated DWARF parser

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

Fixed Issues

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

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

11.0 版本发布

Updated GDB version

• CUDA-GDB已从GDB/7.12升级至GDB/8.2版本。

Support for SM8.0

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

Support for Bfloat16

• 已添加对Bfloat16 (__nv_bfloat16)类型的支持。

MIG support

• CUDA-GDB支持MIG。每个MIG实例上可以运行独立的调试会话。如需使用多个调试器，请参考Multiple Debuggers。

Mac support

• 不再支持在macOS上进行调试。不过，macOS仍可作为主机系统使用（即CUDA-GDB在macOS下运行，通过cuda-gdbserver调试远程目标）。macOS版CUDA-GDB的下载地址如下：Download Here

10.1 版本发布

Enhanced debugging with only linenumber information

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

10.0 版本发布

Turing Uniform Register Support

• 新增支持检查和修改图灵GPU上的统一寄存器。

9.2 版本发布

User induced core dump support

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

9.1 版本发布

Volta-MPS core dump support

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

Lightweight GPU core dump support

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

7.0 版本发布

GPU core dump support

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

6.5 版本发布

CUDA Fortran Support

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

GDB 7.6.2 Code Base

• CUDA-GDB的代码库已升级至GDB 7.6.2版本。

6.0 版本发布

Unified Memory Support

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

GDB 7.6 Code Base

• CUDA-GDB的代码基础从GDB 7.2升级到了GDB 7.6。

Android Support

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

Single-Stepping Optimizations

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

Faster Remote Debugging

• 我们投入了大量努力来显著提升远程调试的速度，最高可提升两个数量级。这些优化同时也加快了本地调试的速度。

Kernel Entry Breakpoints

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

Precise Error Attribution

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

Live Range Optimizations

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

Event Notifications

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

• 新增了内核事件详细程度选项：set cuda kernel_events、set cuda kernel_events_depth。同时set cuda defer_kernel_launch_notifications已被弃用且不再有效。

5.5 版本发布

Kernel Launch Trace

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

Single-GPU Debugging (BETA)

• CUDA-GDB现在可以在渲染桌面GUI的同一GPU上调试CUDA应用程序。此功能还支持调试长时间运行或无限制的CUDA内核，否则会遇到启动超时问题。此外，多个CUDA-GDB会话可以调试在同一GPU上进行上下文切换的CUDA应用程序。此功能在Linux系统上适用于SM3.5设备。有关启用此功能的信息，请参阅桌面管理器运行时的单GPU调试和多调试器。

Remote GPU Debugging

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

5.0 版本发布

Dynamic Parallelism Support

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

Attach/Detach

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

Attach on exception

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

API Error Reporting

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

Inlined Subroutine Support

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

4.2 版本发布

Kepler Support

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

4.1 版本发布

Source Base Upgraded to GDB 7.2

• 在此之前，CUDA-GDB在Linux上基于GDB 6.6版本，在Darwin（苹果分支）上基于GDB 6.3.5版本。现在，两个平台的CUDA-GDB都采用了相同的7.2源代码基础。

• 现在CUDA-GDB支持更新版本的GCC（已测试至GCC 4.5），对DWARF3调试信息提供更好的支持，并增强了C++调试功能。

Simultaneous Sessions Support

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

New Autostep Command

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

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

• 手动单步调试程序是一个缓慢且繁琐的过程。因此，"autostep"功能通过允许用户指定可能发生异常的代码段来辅助用户。当程序运行时，这些代码段会自动进行单步执行，并精确报告在这些代码段内发生的任何异常。

• 在CUDA-GDB中输入'help autostep'可查看该命令的语法和用法。

Multiple Context Support

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

Device Assertions Support

• 与工具包4.1版本一同发布的R285驱动程序支持设备断言功能。CUDA_GDB能够识别断言调用，并在触发断言时暂停应用程序执行。此时可以像往常一样检查变量和内存。如果需要，应用程序也可以继续执行跳过断言。使用'set cuda hide_internal_frames'选项可显示/隐藏系统调用帧（默认隐藏）。

Temporary Directory

• 默认情况下，调试器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进行调试，在编译应用程序时必须向NVCC传递-g -G选项对；例如，

nvcc -g -G foo.cu -o foo

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

• 强制使用-O0编译选项，仅保留极其有限的死代码消除和寄存器溢出优化。

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

注意

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

为了编译您的CUDA Fortran代码并包含CUDA-GDB正常工作所需的调试信息，必须使用PGI CUDA Fortran编译器pgfortran，并加上-g选项。此外，为了便于调试和与未来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 cubins生成代码。对于更新的GPU，内核会在运行时根据目标GPU的架构从PTX重新编译。针对特定虚拟架构进行编译可以确保应用程序在性能权衡的前提下，能够兼容该架构之后的所有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

For additional information, please consult the compiler documentation at https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#extended-notation

3.3. 使用调试器

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

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

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

此外，对于计算能力低于6.0的设备，可以使用软件抢占功能在同一块运行桌面图形界面的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是从info cuda kernels获取的内核ID。

注意

同一个内核可以同时被不同的上下文和设备加载使用。当通过名称或文件名加行号的方式在此类内核中设置断点时，该断点将随机解析到该内核的某一个实例上。使用运行时API时，无法控制断点具体解析到哪个实例。而使用驱动API时，用户可以通过在模块加载后立即设置断点来控制断点解析到哪个实例。

3.3.3. 远程调试

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

按照这种方式设置远程调试需要两个步骤：

在远程主机上启动cuda-gdbserver

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

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

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

  $ cuda-gdbserver :1234 app_invocation
  

  其中1234是cuda-gdbserver将监听的TCP端口号，用于接收来自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环境变量设置为目标根文件系统的位置。如果在运行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命令类似，CUDA命令的内置帮助可以通过在cuda-gdb命令行中使用help命令访问：

(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命令一样，所有CUDA命令都可以通过按TAB键自动补全。

CUDA命令也可以通过apropos命令进行查询。

4.3. 初始化文件

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

4.4. 图形用户界面集成

Emacs

CUDA-GDB 可与 Emacs 和 XEmacs 中的 GUD 协同工作。除了指向正确的二进制文件外，无需额外步骤。

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

确保Emacs/XEmacs的$PATH中包含cuda-gdb。

DDD

CUDA-GDB可与DDD配合使用。要通过DDD使用CUDA-GDB，请使用以下命令启动DDD：

ddd --debugger cuda-gdb

cuda-gdb 必须位于您的 $PATH 环境变量中。

4.5. GPU核心转储支持

有两种方式可以为CUDA应用程序配置核心转储选项。在应用程序环境中设置环境变量，或通过应用程序使用CUDA Driver 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核心转储时改变默认生成行为。可以向该环境变量提供多个标志，并用,分隔。这些标志可用于实现诸如减小生成的GPU核心转储大小或其他与默认行为不同的期望行为。下表列出了每个标志及其存在时的行为。

GPU核心转储 CUDA_COREDUMP_GENERATION_FLAGS

环境变量标志	描述
skip_nonrelocated_elf_images	禁用将未重定位的ELF镜像副本包含在GPU核心转储中。仅保留重定位后的镜像。
skip_global_memory	禁用GPU全局内存和常量存储区内存段的转储。
skip_shared_memory	禁用GPU共享内存段的转储。
skip_local_memory	禁用GPU本地内存段的转储。
skip_abort	在GPU核心转储生成过程结束时禁用调用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"。

核心转储生成的限制和注意事项

核心转储支持存在以下限制：

• 对于Windows WDDM，GPU核心转储仅支持计算能力6.0或更高的GPU。Windows TCC支持所有受支持计算能力上的GPU核心转储。

• 对于运行在SLI模式下的GPU，Windows Subsystem for Linux不支持GPU核心转储。支持多GPU设置，但无法在驱动程序控制面板中启用SLI模式。

• 仅当启用硬件调度模式时，Windows Subsystem for Linux才支持GPU核心转储。

• 由于在Windows子系统Linux中执行Linux程序时存在NTFS文件系统的限制，仅当被执行的二进制文件位于Linux文件系统上时，才支持GPU核心转储功能。

• 不支持在Windows Subsystem for Linux环境下使用CUDA_ENABLE_USER_TRIGGERED_COREDUMP选项生成GPU核心转储。

• 在QNX平台上，目前不支持使用CUDA_ENABLE_CPU_COREDUMP_ON_EXCEPTION生成CPU核心转储。

• NVIDIA CMP产品线不支持GPU核心转储功能。

• 每个上下文的崩溃转储只能在计算能力6.0或更高的GPU上启用。计算能力低于6.0的GPU在使用崩溃转储属性控制API时将返回CUDA_ERROR_NOT_SUPPORTED。

• 如果MPS客户端触发了核心转储(core dump)，运行在同一MPS服务器上的其他所有客户端都会出现故障。这些间接受影响的客户端如果启用了核心转储生成功能，同样会产生核心转储文件。

• 当其他开发者工具（包括CUDA-GDB）正在与应用程序交互时，不支持GPU核心转储功能。除非明确文档说明为支持的用例（例如generate-cuda-core-file命令）。

• 在异常时生成核心转储文件时，如果内核在识别到异常之前退出，可能导致无法生成核心文件。有关如何解决此问题的策略，请参阅GPU错误报告中的说明。

注意

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

注意

Starting from CUDA 11.6, the compute-sanitizer tool can generate a GPU core dump when an error is detected by using the --generate-coredump yes option. Once the core dump is generated, the target application will abort. See the compute-sanitizer documentation for more information: https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html#coredump

注意

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

注意

NVIDIA vGPU平台必须显式启用调试支持才能生成GPU核心转储文件。有关如何在vGPU上启用调试的信息，请参阅Virtual GPU Software User Guide。

注意

NVIDIA Jetson和Drive Tegra设备必须显式启用调试支持才能生成GPU核心转储文件。请参阅在Jetson和Drive Tegra设备上使用CUDA-GDB调试器章节。

注意

在运行QNX的NVIDIA Drive Tegra设备上生成核心转储时，生成CPU核心转储可能会导致系统挂起。如果遇到挂起情况，请将CUDA_ENABLE_CPU_COREDUMP_EXCEPTION设置为0。

注意

如果使用OptiX/RTCore构建的程序运行时未生成核心转储文件，请尝试将环境变量OPTIX_FORCE_DEPRECATED_LAUNCHER设置为1。具体请参阅调试OptiX/RTCore应用程序章节。

注意

在Windows和Windows Subsystem for Linux (WSL)上生成核心转储文件需要通过将注册表键>HKEY_LOCAL_MACHINE\SOFTWARE\NVIDIA Corporation\GPUDebugger\EnableInterface设置为(DWORD) 1来启用调试接口。

注意

GPU核心转储仅在启用开发者工具模式的机密计算模式下运行的GPU上受支持。有关如何启用该模式的更多详细信息，请参阅机密计算部署指南。

GPU核心转储文件的命名

默认情况下，GPU核心转储文件会在当前工作目录中生成。其命名为core_TIME_HOSTNAME_PID.nvcudmp，其中TIME表示自纪元以来的秒数，HOSTNAME是运行CUDA应用程序的主机名，PID则是CUDA应用程序的进程标识符。

CUDA_COREDUMP_FILE 环境变量可用于定义模板，该模板用于更改GPU核心转储文件的名称。模板可以是绝对路径，也可以是相对于当前工作目录的相对路径。模板可以包含%说明符，当创建GPU核心转储时，这些说明符会被以下模式替换：

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

例如，将 CUDA_COREDUMP_FILE 设置为：

export CUDA_COREDUMP_FILE=newName.%h.%p

这将导致GPU核心转储文件被写入当前工作目录下的newName.myhost.1234。其中myhost和1234会分别被替换为实际的主机名和进程ID。

将 CUDA_COREDUMP_FILE 设置为：

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

将导致GPU核心转储文件按照与上述示例相同的命名逻辑写入用户的主目录。

如果CUDA_COREDUMP_FILE指向一个现有的FIFO类型文件（例如命名管道），核心转储将被流式传输到该文件。

核心转储文件可以通过CUDA_COREDUMP_FILE管道传输到shell命令，格式如下：

export CUDA_COREDUMP_FILE='| cmd > file'

例如，要将核心转储通过管道传输到 gzip 使用：

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

注意

当传输核心转储时，% 说明符将不会被识别。

启用用户触发的GPU核心转储生成

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

将环境变量CUDA_ENABLE_USER_TRIGGERED_COREDUMP设置为1以启用用户触发的GPU核心转储功能。该选项默认处于禁用状态。设置此环境变量将为每个后续运行的CUDA进程打开一个通信管道。用户只需向该管道写入数据即可触发GPU核心转储。

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

可以使用cuda-gdb的generate-core-file或gcore命令手动生成GPU核心转储文件。可以选择性地传入文件名，详情请参阅help generate-core-file。提供多个标志位用于控制保存内容以及是否同时生成CPU核心转储文件。

显示核心转储生成进度

默认情况下，当应用程序崩溃并生成GPU核心转储时，在完全生成之前应用程序可能会显示为无响应或冻结状态。

将环境变量CUDA_COREDUMP_SHOW_PROGRESS设置为1，以便将核心转储生成进度信息打印到stderr。这可用于确定核心转储生成的进度：

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核心转储生成

Driver API为所有环境变量提供了等效的设置选项，并新增了能够针对每个上下文而非全局设置不同核心转储配置的功能。该API可直接在应用程序中调用。使用cuCoredumpGetAttributeGlobal和cuCoredumpSetAttributeGlobal来获取或设置全局属性。使用cuCoredumpGetAttribute和cuCoredumpSetAttribute来获取或设置上下文级属性。更多信息请参阅Coredump Attributes Control API手册。

下表列出了环境变量及对应的CUcoredumpSettings标志，这些变量和标志可用于通过核心转储属性控制API来管理核心转储设置。

注意

CU_COREDUMP_ENABLE_USER_TRIGGER 设置只能在驱动API中全局配置，且必须先设置好CU_COREDUMP_PIPE（如需使用）才能启用用户触发的核心转储功能。

GPU核心转储配置参数

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

在cuda-gdb中检查GPU和GPU+CPU核心转储

使用以下命令将GPU核心转储加载到调试器中

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

  这将打开核心转储文件并打印程序执行期间遇到的异常。然后，可以发出标准的cuda-gdb命令来进一步调查应用程序在终止时的设备状态。

使用以下命令将CPU和GPU核心转储加载到调试器中

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

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

注意

核心转储检查不需要在系统上安装GPU

5. 内核聚焦

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

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

5.1. 软件坐标与硬件坐标

设备线程属于一个块，而块又属于一个内核。线程、块和内核是关注点的软件坐标。设备线程在通道上运行。通道属于一个线程束，线程束属于一个流式多处理器(SM)，而SM又属于一个设备。通道、线程束、SM和设备是关注点的硬件坐标。只要保持一致性，软件坐标和硬件坐标可以互换使用或同时使用。

有时会使用另一种软件坐标：网格(grid)。网格与内核(kernel)的区别在于作用范围。网格ID在单个GPU内是唯一的，而内核ID在所有GPU之间都是唯一的。因此，内核与(网格,设备)元组之间存在一一对应的映射关系。

注意

如果启用了软件抢占功能（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的值。当该值为开启状态（默认）时，聚焦的warp将持续单步执行，直到聚焦的CUDA线程变为活动状态。当该值为关闭状态时，聚焦的warp将被执行单步，且聚焦的CUDA线程将更改为该warp中最近的活动通道。

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

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

只要设备函数没有被内联，你就可以单步进入、单步跳过或单步跳出这些函数。要强制编译器不对函数进行内联，必须在函数声明中添加__noinline__关键字。

在设备上执行的异步SASS指令（如warpgroup指令）在之前的程序计数器（PC）处不能保证已完成。

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

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

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

• 当单步执行设备网格启动至完成时，焦点会自动切换回CPU。必须使用cuda kernel焦点切换命令才能切换到另一个目标网格（如果该网格仍驻留在内存中）。

注意

无法单步进入设备启动调用（也无法进入该调用启动的例程）。

7. 断点与监视点

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

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

设置断点时，当所有驻留的GPU线程到达对应的程序计数器位置时，会强制它们在此处停止。

当一个线程触发断点时，无法保证其他线程会同时触发该断点。因此同一个断点可能会被多次触发，用户必须仔细检查实际触发断点的线程。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)

也可以使用函数的混淆名称。要查找函数的混淆名称，可以使用以下命令：

(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. CUDA信息命令

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

devices

关于所有设备的信息

sms

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

warps

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

lanes

关于当前warp中所有活动通道的信息

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 设备

该命令会枚举系统中所有按设备索引排序的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. cuda sms信息

该命令显示设备的所有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. cuda线程束信息

该命令可让您深入一层，打印出当前关注的SM中所有线程束的信息。此命令支持过滤器，默认值为device current sm current warp all。该命令可用于显示每个线程束正在执行哪个块。

(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. cuda通道信息

该命令显示当前关注的warp中的所有通道(线程)。此命令支持过滤器，默认值为device current sm current warp current lane all。在下面的示例中，您可以看到所有通道都位于相同的物理PC位置。该命令可用于显示每个通道执行的是哪个线程。

(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. cuda内核信息

该命令显示当前聚焦GPU上所有活跃的内核。它会打印每个内核的SM掩码、内核ID和网格ID，以及相关的维度和参数。内核ID在所有GPU中是唯一的，而网格ID在每个GPU内是唯一的。Parent列显示父级网格的内核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上启动的网格。具有负网格ID的内核是从GPU启动的，而具有正网格ID的内核是从CPU启动的。

8.3.6. CUDA块信息

该命令显示当前焦点内核中所有活跃或正在运行的块。结果按内核分组显示。此命令支持过滤器，默认值为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 线程

该命令显示应用程序当前活跃的CUDA块和线程，以及这些块中的线程总数。同时还会显示虚拟程序计数器(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
 ...

注意

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

该命令还支持breakpoint all和breakpoint breakpoint_number作为筛选条件。前者显示命中用户设置的所有CUDA断点的线程，后者显示命中指定breakpoint_number号CUDA断点的线程。

(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. cuda启动追踪信息

此命令显示焦点内核的内核启动追踪。追踪中的第一个元素即为焦点内核。下一个元素是启动该内核的内核。追踪将持续进行，直到没有父内核为止。这种情况下，内核是由CPU启动的。

对于跟踪中的每个内核，该命令会打印内核在跟踪中的层级、内核ID、设备ID、网格ID、状态、内核维度、内核名称以及内核参数。

(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上运行的内核将显示为Pending状态。当前正在GPU上运行的内核会被标记为Active。等待再次激活的内核将显示为Sleeping。当内核终止时，它会被标记为Terminated。在少数情况下，当调试器无法确定内核是挂起还是终止时，状态会被设置为Undetermined。

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

注意

使用set cuda software_preemption on时，不会报告任何内核处于活动状态。

8.3.9. cuda启动子进程信息

该命令显示当前聚焦内核启动的所有未终止内核列表。对于每个内核，会显示内核ID、设备ID、网格ID、内核维度、内核名称以及内核参数。

(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. CUDA上下文信息

该命令会枚举所有GPU上正在运行的CUDA上下文。*标记表示当前处于焦点状态的上下文。此命令可显示某个上下文当前是否在设备上处于活动状态。

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

8.3.11. cuda托管内存信息

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

(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和错误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访问，例如：

(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. 常量存储体

分配在GPU内存常量地址空间中的内存驻留在称为常量存储体的二维数组中。 常量存储体标记为c[X][Y]，其中X是存储体编号，Y是偏移量。 特定存储体/偏移量对的内存地址可通过便捷函数$_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会向用户通知内核事件和上下文事件。在CUDA-GDB中，kernel指的是在GPU上执行的设备代码，而context指的是GPU上为内核分配的虚拟地址空间。您可以启用CUDA上下文和内核事件的输出来查看活动上下文和内核的流程。默认情况下，仅显示上下文事件消息。

9.1. 上下文事件

每当应用程序创建、推送、弹出或销毁CUDA上下文时，CUDA-GDB可以选择性地显示通知消息。该消息包含上下文ID及其所属的设备ID。

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

默认情况下，上下文事件通知处于禁用状态。上下文事件通知策略通过context_events选项进行控制。

• (cuda-gdb) set cuda context_events off
  

  CUDA-GDB默认不显示上下文事件通知消息。

• (cuda-gdb) set cuda context_events on
  

  CUDA-GDB将显示上下文事件通知消息。

9.2. 内核事件

每当CUDA-GDB感知到CUDA内核的启动或终止时，都会显示通知消息。该消息包含内核ID、内核名称以及该内核所属的设备。

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

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

• (cuda-gdb) 设置 cuda 内核事件为无
  

可能的选项包括：

none

无内核、应用程序或系统（默认）

application

由用户应用程序启动的内核

system

由驱动程序启动的任何内核，例如memset

all

任何内核、应用程序和系统

• (cuda-gdb) set cuda kernel_events_depth 0
  

  控制内核事件通知显示的最大深度。值为0表示没有限制，将显示所有内核通知。值为1表示调试器仅显示从CPU启动的内核事件通知（默认值）。

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上报告的新错误。

注意

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

注意

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

• 对于Volta及以上架构，使用-G编译应用程序。更多信息请参阅Compiling the Application。

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

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

CUDA异常代码

异常代码	错误精度	错误范围	描述
CUDA_EXCEPTION_0 : "设备未知异常"	未知	GPU上的全局错误	这是一个由应用程序引起的全局GPU错误，与下列列出的任何错误代码都不匹配。这种情况应该很少发生。可能的原因包括Device Hardware Stack溢出，或者内核在即将终止时生成了异常。
CUDA_EXCEPTION_1 : "已弃用"	已弃用	已弃用	此异常已弃用，应视为CUDA_EXCEPTION_0处理。
CUDA_EXCEPTION_2 : "线程用户栈溢出"	精确	每线程错误	当线程超出其栈内存限制时会发生此错误。
CUDA_EXCEPTION_3 : "设备硬件堆栈溢出"	精确	GPU上的全局错误	当应用程序触发全局硬件堆栈溢出时会发生此错误。此错误的主要原因是存在函数调用时出现大量分支发散。
CUDA_EXCEPTION_4 : "Warp非法指令"	精确	Warp错误	当warp中的任何线程执行了非法指令时会发生此错误。
CUDA_EXCEPTION_5 : "Warp越界地址"	精确	Warp错误	当一个warp中的任何线程访问超出本地或共享内存区域有效范围的地址时，会发生此错误。
CUDA_EXCEPTION_6 : "线程束地址未对齐"	精确错误	线程束错误	当线程束中的任意线程访问本地内存或共享内存段中未正确对齐的地址时，会触发此异常。
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 Assert"	精确	每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 : "集群目标块不存在"	不精确	每个Cuda集群	当块内的任何线程访问属于集群的有效块范围之外的另一个块时，会发生此情况。
CUDA_EXCEPTION_18 : "集群地址越界"	不精确	每个Cuda集群	当块内任何线程访问超出该集群所属共享内存区域有效范围的地址时，会触发此异常。

10.3. 自动步进

Autostep是一个命令，用于将CUDA异常的精度提升到具体的执行通道和指令级别，这在其他情况下是无法实现的。

在正常执行情况下，异常可能在发生后的若干条指令后才被报告，或者除非是通道错误，否则可能无法确定发生异常的确切线程。然而，如果在异常发生时程序正处于单步执行状态，就能精确定位异常源头。手动单步调试是一个缓慢且繁琐的过程：单步执行耗时远超过正常执行，且用户必须逐个线程束进行单步调试。

Autostep通过允许用户指定他们认为可能发生异常的代码段来辅助用户，这些代码段在程序运行时会被自动且透明地单步执行。程序的其余部分正常执行，以最小化单步执行导致的减速。如果异常发生在这些代码段内，将报告异常的确切来源。因此，通过使用autostep，可以快速且更轻松地找到发生异常的精确指令和线程。

注意

autostep 命令是通过断点实现的。如果在当前程序计数器(PC)上设置了自动步进(autostep)，该自动步进功能将不会生效，直到下次再次遇到该程序计数器时才会被评估。

自动步进使用说明

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

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

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

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

• 在自动步进过程中调用的函数将被跳过。

• 在出现分歧的情况下，自动步进窗口的长度由每个warp中第一个活动线程执行的指令行数或指令数量决定。
  分歧线程也会单步执行，但它们执行的指令不计入自动步进窗口的长度。

• 如果在自动步进窗口内触发断点，当程序恢复运行时，命中断点的线程束(warp)将不会继续自动步进。但其他线程束仍可能保持自动步进状态。

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

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

如果在内存错误位置之前设置了自动步进（autostep）但未触发内存错误，则可能是所选窗口太小所致。这种情况可能是由于自动步进位置地址与触发内存错误的指令之间存在函数调用所导致。此时，要么增大窗口大小以确保包含错误指令，要么将自动步进位置移动到更接近错误指令执行时间的指令处。

相关命令

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

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

• 示例：自动步进

• 示例：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且编译时不需要额外的编译器标志。另请参阅Debug Compilation

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. 此时可以输入命令来推进执行或打印程序状态。在本教程中，我们将继续执行直到设备内核启动。

   (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设备内核。调试器会打印当前聚焦的CUDA线程。

6. 使用info cuda threads命令查看当前聚焦的CUDA线程，并在主机线程和CUDA线程之间切换：

   (cuda-gdb) info cuda threads
     块索引   线程索引   目标块索引 线程索引 数量            虚拟PC
   文件名      行号
   
   内核0
   * (0,0,0)   (0,0,0)   (0,0,0)  (255,0,0)   256 0x0000000000866400 bitreverse.cu     9
   (cuda-gdb) thread
   [当前线程是1 (进程16738)]
   (cuda-gdb) thread 1
   [切换到线程1 (进程16738)]
   #0  0x000019d5 在 main () 位于 bitreverse.cu:34
   34    bitreverse<<<1, N, N*sizeof(int)>>>(d);
   (cuda-gdb) backtrace
   #0  0x000019d5 在 main () 位于 bitreverse.cu:34
   (cuda-gdb) info cuda kernels
   内核 设备 网格   SMs掩码 网格维度 块维度        名称 参数
        0   0    1 0x00000001 (1,1,1) (256,1,1) bitreverse data=0x110000
   (cuda-gdb) cuda kernel 0
   [将焦点切换到CUDA内核0，网格1，块(0,0,0)，线程(0,0,0)，设备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) 位于 bitreverse.cu:9
   

7. 通过打印块和线程索引来验证此信息：

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

8. 也可以打印网格和块的维度：

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

9. 推进内核执行并验证一些数据：

   (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.
    
    程序正常退出。
    (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 }

在这个小例子中，我们有一个指向整数的指针数组，我们想对这些整数进行一些操作。然而，假设其中一个指针如第38行所示是NULL。当我们尝试访问与块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错误。然而，报告的线程是块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错误，因此我们知道抛出异常的线程必须与块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_Comm_rank函数基于MPI等级来条件化自旋循环。

对于计算能力低于6.0的设备，Multiple Debuggers中描述的软件抢占解决方案不适用于MPI应用程序。对于这些GPU，请确保每个MPI进程对应唯一的GPU。

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

12. 技巧与窍门

本节作为高级设置和各种技巧的参考指南，供CUDA-GDB用户使用，这些内容在其他文档中未被记录。

12.1. 设置cuda break_on_launch

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

(cuda-gdb) set cuda break_on_launch application

可能的选项包括：

none

无内核、应用程序或系统（默认）

application

由用户应用程序启动的内核

system

由驱动程序启动的任何内核，例如memset

all

任何内核、应用程序和系统

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

12.2. 设置 cuda 启动阻塞模式

启用后，内核启动将变为同步模式，就像环境变量CUDA_LAUNCH_BLOCKING被设置为1一样。一旦进入阻塞状态，内核启动实际上会被序列化，可能更易于调试。

• (cuda-gdb) set cuda launch_blocking off
  

  内核启动将根据应用程序的要求同步或异步执行。这是默认设置。

• (cuda-gdb) set cuda launch_blocking on
  

  内核启动将变为同步模式。如果应用程序已经启动，该更改只会在当前会话结束后生效。

12.3. 设置cuda通知

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

• (cuda-gdb) set cuda notify youngest
  

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

• (cuda-gdb) set cuda notify random
  

  任意一个主机线程将收到通知信号。

12.4. 设置cuda ptx缓存

在访问变量值之前，调试器会检查该变量在当前程序计数器(PC)位置是否存活。在CUDA设备上，变量可能不会一直保持存活状态，并会被报告为"已优化移除"。

CUDA-GDB提供了一种绕过此限制的选项，通过缓存PTX寄存器级别的变量值。每个源变量会被编译成一个PTX寄存器，随后映射到一个或多个硬件寄存器。利用编译器发出的调试信息，调试器可以根据变量先前映射的最新硬件寄存器来缓存PTX寄存器的值。

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

• (cuda-gdb) set cuda ptx_cache off
  

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

• (cuda-gdb) set cuda ptx_cache on
  

  调试器将尽可能使用缓存值。此设置为默认设置且始终安全。

12.5. 设置CUDA单步调试优化

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

• (cuda-gdb) set cuda single_stepping_optimizations off
  

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

• (cuda-gdb) set cuda single_stepping_optimizations on
  

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

12.6. 设置cuda线程选择

当调试器需要选择一个活动线程进行聚焦时，该决策由启发式算法指导。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. 设置CUDA值外推

在访问变量值之前，调试器会检查该变量在当前程序计数器(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运行命令中：

--cap-add=SYS_PTRACE

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

随着CTK 11.8版本的发布，Linux平台引入了一个名为统一调试器(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子系统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 Toolkit 发行说明。

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

适用于Linux的Windows子系统(WSL)

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

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

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

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

以下是当前版本在支持的操作系统上的已知问题及其解决方法。

Python未初始化

出现此问题的原因是机器上缺少Python 3.x库，安装该库即可解决。当PATH环境变量中默认python3解释器安装的libpython主次版本不匹配时也会引发此问题。必须确保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应用程序上使用cuda-gdb进行附加调试。

• 不支持在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设备上，GPU调试功能仅在调试器以root权限运行时才能正常工作。若要在非root权限下运行调试器，需要修改devfs节点权限设置。

• 调试器可能会漏报一个引发的陷阱(__trap())，如果该陷阱是设备从断点恢复后执行的下一条指令。

• 在恢复执行过程中，如果先前空闲的SM上启动了新的warp，调试器可能会漏报断点或异常。

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

• 调试器不支持访问通过CUDA IPC API从其他进程导入的共享内存分配。调试器尝试访问这些共享内存分配时，将返回错误提示：不支持访问通过IPC共享的内存分配。

• 除非将OPTIX_FORCE_DEPRECATED_LAUNCHER设置为1，否则break_on_launch在OptiX/RTCore程序中不会生效。

• 在QNX系统上，需要将QNX_TARGET环境变量指向目标根文件系统。如果该变量不可用，请将环境变量设置为空字符串。

16. 通知

16.1. 通知

本文档仅供信息参考之用，不应视为对产品功能、状态或质量的保证。NVIDIA公司（“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对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。

16.2. OpenCL

OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。

16.3. 商标

NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。