Debugging CUDA: An Overview of CUDA Correctness Tools

Steve Ulrich & Aurelien Chartier | GTC March 2023

1. 调试器工具
- 1.1 调试器工具概览
- 1.2 为调试准备代码 (Getting your Code Ready for Debugging
2. CUDA GDB
3. Nsight Visual Studio Edition
4. Nsight Visual Studio Code Edition
- 4.1 安装与配置
5. Nsight Eclipse Edition
- 5.1 安装与配置
- 5.2 Java 支持
6. CUDA Debugger API
7. Compute Sanitizer
8. 如何获取工具
9. GTC 相关的开发者工具资源

1. 调试器工具

本节将概述多种NVIDIA提供的调试器工具。

1.1 调试器工具概览

NVIDIA提供的调试工具可分为三类：

命令行工具 (Command Line Tools)
- CUDA GDB (适用于 Linux)
- Sanitizer
IDE 集成工具 (IDE Tools)
- Nsight Eclipse Edition (适用于 Linux)
- Nsight Visual Studio Edition (适用于 Windows)
- Nsight Visual Studio Code Edition
开发者库 (Developer Libraries)
- CUDA Debugger API (适用于 Linux)
- Sanitizer API

1.2 为调试准备代码 (Getting your Code Ready for Debugging)

调试器的效果取决于编译器提供的元数据 (metadata)。
为调试而编译:
- -g - 编译 CPU 代码以进行调试。
- -G - 编译 GPU 代码以进行调试。
副作用 (Side effects):
- 编译器会将元数据插入到生成的可执行文件中，以指导调试器工作，包括：
  - 局部变量、参数、静态变量、全局变量等的位置。
  - 如何从当前函数回溯调用堆栈，直至调用图的根节点。
- 性能会显著下降。
- 反汇编代码会变得"平淡无奇" (unimpressive) - 包含冗余的加载/存储指令等。

2. CUDA GDB

2.1 概述 (Overview)

CUDA GDB 是一个强大的命令行调试工具。

基于熟悉的 GDB 调试器构建
- 易于使用：已熟悉 GDB 的用户可以轻松上手。
- GPU 调试提供与 GDB 类似的逻辑体验。
- 支持 C/C++/Fortran。
- 在主机 (CPU) 和设备 (GPU) 调试之间提供无缝体验。
- 支持 CUDA/OptiX/etc 的源码级设备代码。
- 支持 SASS 反汇编。
- 包含 CUDA GDB 特有的各种命令扩展。
交互式 CLI 工具
- 提供对 CUDA kernels 的调试。
- 处理 CUDA 运行时错误。
- 调试导致不正确结果的逻辑错误。
事后调试 (Post-mortem debugging) 支持 corefiles
- 通过环境变量捕获 coredump。
基于 GDB 12.1
完整源代码可在 https://github.com/NVIDIA/cuda-gdb 获取

2.2 架构支持 (Architecture Support)

Linux
- 支持的发行版包括：CentOS / Debian / Fedora / KylinOS / OpenSUSE / RHEL / SLES / Ubuntu
- 支持 x86 和 aarch64 架构。
Windows (WSL2)
Mac
- 仅支持主机端调试 (无编译器)。
- 通过 "target" 连接到远程 CUDA GDBSERVER。
- 可通过单独下载获取 (非标准 CUDA Toolkit 的一部分)。

2.3 基础操作 (Basics)

两种获取控制权的方式
- run: 启动并调试应用

$ cuda-gdb --quiet my_application
Reading symbols from my_application...
(cuda-gdb) run

* **attach**: 附加到已在运行的进程

$ cuda-gdb --quiet
(cuda-gdb) attach 261230

退出调试器
- 使用 quit 命令。
- 被调试的应用将被终止。
- 已附加的应用将被分离。
恢复应用执行
- 使用 (cuda-gdb) continue 命令。
- 同时恢复主机和设备线程。
中断执行
- 应用正在执行时，按 Ctrl-C 可暂停主机和设备线程。
- 此时不会显示 (cuda-gdb) 提示符。

使用 info cuda 命令查询 CUDA GPU 活动
- 该命令可打印当前 CUDA 活动的信息，提供多种可用选项。

Page 8 a `cuda-gdb` command output showing the `help info cuda` command.

CUDA 线程焦点由 cuda 命令控制
- 将焦点设置到单个 CUDA 线程。
- 某些命令仅适用于当前焦点的线程，例如：
  - 打印局部或共享变量。
  - 打印寄存器。
  - 打印堆栈内容。
示例
- 将焦点设置到指定的 CUDA 线程：

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

* 基于 block 和 thread 设置焦点：

(cuda-gdb) cuda kernel 0 block 2 thread 6
[Switching focus to CUDA kernel 0, grid 1, block (2,0,0), thread (6,0,0), device 0, sm 4, warp 0, lane 6]

* 基于 kernel, block, dim3 thread 设置焦点：

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

disassemble (反汇编)
- 查看 SASS 指令。
- 当前 PC (程序计数器) 前缀为 =>。
- 触发异常的指令前缀为 *>。
- 如果指令既是当前 PC 又是异常触发点，则前缀为 *=>。
示例输出：

(cuda-gdb) disas $pc,+32
Dump of assembler code from 0x7fffc385b4b0 to 0x7fffc385b4d0:
=> 0x00007fffc385b4b0 <_Z16exception_kernelPv11exception_t+3504>:	ERRBAR
   0x00007fffc385b4c0 <_Z16exception_kernelPv11exception_t+3520>:	EXIT
End of assembler dump.

Dump of assembler code from 0x7fffc385ab20 to 0x7fffc385ab60:
*> 0x00007fffc385ab20 <_Z16exception_kernelPv11exception_t+1056>:	ST.E.U8.STRONG.SYS DESC[UR4][R6.64], R5
   0x00007fffc385ab30 <_Z16exception_kernelPv11exception_t+1072>:	BRA 0xad0
   0x00007fffc385ab40 <_Z16exception_kernelPv11exception_t+1088>:	PRMT R5, R5, 0x7610, R5
   0x00007fffc385ab50 <_Z16exception_kernelPv11exception_t+1104>:	MOV R6, c[0x0][0xc]
End of assembler dump.

* **注意**：不支持 PTX 反汇编。

2.4 Coredumps

GPU coredump 支持
- 默认禁用。
- 设置环境变量 CUDA_ENABLE_COREDUMP_ON_EXCEPTION 为 1 来启用。
- 当 GPU 发生异常时会生成 coredump 文件。

$ ./memexceptions 1
SM version: 86, Min version: 35, Max version: 999
Aborted (core dumped)
$ ls | grep core
core_1669651659_agontarek-dt_612954.nvcudmp

GPU coredump 文件名
- 默认格式为 core_%t_%h_%p.nvcudmp
- %t 是自 Epoch 以来的秒数。
- %h 是运行 CUDA 应用的系统主机名。
- %p 是 CUDA 应用的进程标识符。
- 默认写入应用的 $PWD 目录。
用户可通过 CUDA_COREDUMP_FILE 环境变量自定义
- 识别 %t, %h, %p 等说明符。

$ export CUDA_COREDUMP_FILE="/lus/grand/projects/alcf_training/$USER/core.gpu.%h.%p"

2.5 多 GPU 调试 (Multi-GPU Debugging)

支持拥有多个 GPU 的系统。
断点会停止所有正在运行 CUDA 的 GPU。
使用 "info cuda kernels" 列出活动的 kernels。
使用 "cuda kernel <n>" 在 kernels 之间切换。
受影响的 GPU 可通过环境变量 CUDA_VISIBLE_DEVICES 控制。

2.6 Python 支持 (Python Support)

支持 GDB 的 Python 解释器。
基于 Python 3.6m 构建。
通过 dlopen 加载。
注意事项 (Caveats)
- CUDA GDB 遵循 "一次构建，到处运行" 的原则。
- 并非所有发行版都拥有兼容的 Python 库。
- "--disable-python" 选项可以跳过 Python 初始化。

Python 脚本示例：
Page 13 a code block showing a python script being run in `cuda-gdb`.

2.7 WSL2 支持 (WSL2 Support)

可与 Microsoft 的 WSL2 (Windows Subsystem for Linux 2) 配合使用。
支持从 Pascal 到 Ampere 的架构。
对 Ada 和 Hopper 架构的支持即将推出。

3. Nsight Visual Studio Edition

Nsight Visual Studio Edition 是一个与 Microsoft Visual Studio 完全集成的 IDE，为 CUDA C/C++ 开发提供了强大的 GPU 调试功能。

支持 Visual Studio 2017, 2019 和 2022
CUDA C/C++ GPU 调试
源码关联的汇编级调试 (SASS / PTX / SASS+PTX)
支持 CUDA C/C++ 代码的数据断点
在 Locals, Watch 和 Conditionals 窗口中支持表达式

Nsight Visual Studio Edition 提供了一个集成的调试环境，用户可以在 Visual Studio 中直接调试 CUDA 代码。界面中集成了多个用于 GPU 调试的专用窗口。

Page 16: Nsight Visual Studio Edition 界面展示了GPU寄存器、本地变量、断点和CUDA C++源代码。

上图展示了调试界面中的关键组件：
* GPU 寄存器 (GPU Registers): 显示当前选定线程的寄存器值。
* 本地变量 (Locals): 显示当前作用域内的变量值。
* 断点 (Breakpoints): 管理代码中断点。

3.1 反汇编视图 (Disassembly View)

Nsight VSE 允许开发者查看 CUDA kernel 的底层代码，支持 SASS（设备原生汇编）、PTX（并行线程执行）或两者的混合视图，有助于进行低级别的性能分析和调试。

Page 17: Nsight Visual Studio Edition 的反汇编视图，可以选择查看 SASS、PTX 或两者。

3.2 条件断点 (Conditional Breakpoints)

开发者可以设置条件断点，仅当满足特定条件时（例如，特定线程ID执行时）才触发中断。这对于调试大规模并行代码中的特定线程或问题非常有用。

Page 18: 在 Nsight Visual Studio Edition 中设置条件断点，示例条件为 threadIdx.x == 3 && threadIdx.y == 7。

3.3 Warp 信息 (Warp Info)

Warp Info 窗口提供了对 Warp 级别的深入洞察，显示了 Warp 中所有线程的状态，包括它们所在的 SM（流多处理器）ID、Grid ID、线程 ID 等信息。这有助于理解 Warp 内的执行分歧和线程状态。

Page 19: Nsight Visual Studio Edition 的 Warp Info 窗口，展示了一个 Warp 中所有线程的详细信息。

3.4 GPU 架构支持

Nsight Visual Studio Edition 的调试器后端支持不同的 GPU 架构：
* 使用统一调试器后端 (Unified Debugger backend) 支持 Pascal 及更新的架构。
* 使用传统调试器后端 (Legacy Debugger backend) 支持 Maxwell 架构。

Page 20: Nsight Visual Studio Edition 的 GPU 架构支持。

4. Nsight Visual Studio Code Edition

Nsight Visual Studio Code Edition 是为 CUDA 开发者提供的 Visual Studio Code IDE 扩展。

作为 Microsoft Visual Studio Code 的扩展提供。
构建于 CUDA GDB 之上。
提供 CUDA 语言支持，包括：
- IntelliSense（智能感知）
- 语法高亮
- 问题匹配器 (Problem matcher)
支持 CUDA C/C++ 的 CPU/GPU 联合调试。

Page 21: Nsight Visual Studio Code Edition 及其特性。

下图展示了在 Visual Studio Code 中调试 CUDA 代码的界面，包括变量监视、调用堆栈、断点管理和集成的终端。

Page 22: Nsight Visual Studio Code Edition 调试界面。

4.1 安装与配置

Nsight VS Code Edition 的安装方式灵活：
* 通过扩展安装: 可在 Visual Studio Marketplace 中获取。
* 通过下载安装: 可从 NVIDIA 开发者网站下载：https://developer.nvidia.com/nsight-visual-studio-code-edition
* 安装程序与 CUDA toolkit 安装程序是解耦的（独立）。
* 需要单独安装 CUDA toolkit 以支持编译（通过编译器）和调试（通过 CUDA GDB）。

Page 23: Nsight Visual Studio Code Edition 的安装选项。

5. Nsight Eclipse Edition

Nsight Eclipse Edition 是一个用于 CUDA 开发的 Eclipse IDE。

构建于 CUDA GDB 之上。
功能齐全的 IDE，用于编辑、构建和调试 CUDA 应用程序。
可以将 Nsight Eclipse 插件安装到用户自己的 Eclipse 环境中。
支持的 Eclipse 版本 - (Java 8 / Java 11)。
NVCC 构建集成支持交叉编译。
支持的平台：(x86/L4T/Drive Linux/Drive QNX)。
使用 CUDA GDB 同时调试 CPU 和 GPU 代码。

下图展示了 Nsight Eclipse Edition 的调试界面，包括项目浏览器、源代码编辑器以及显示 CUDA 内置变量的变量视图。

5.1 安装与配置

通过扩展安装（随 CUDA toolkit 提供）。
插件可以在 /usr/local/cuda/nsightee_plugins 目录中找到。

5.2 Java 支持

Java 8 支持
- 经过测试的支持 - Eclipse 4.16
- 版本 4.9 到 4.15 很可能可以工作，但未经测试。
Java 11 支持
- 经过测试的支持 - Eclipse 4.19, 4.24, 4.25
- 版本 4.20 到 4.23 很可能可以工作，但未经测试。

Page 27: Nsight Eclipse Edition 的 Java 版本支持详情。

6. CUDA Debugger API

CUDA Debugger API 旨在支持第三方调试器的集成。

仅限 Linux。
支持 ABI（应用程序二进制接口）。
异常报告。
附加 (Attach) 和分离 (Detach) 进程。
运行时控制。
状态检查。

该 API 是一个底层接口，上层工具如 Nsight Eclipse Edition、Nsight Visual Studio Code、Arm/Linaro Forge 和 Perforce TotalView 通过 CUDA GDB 使用此 API，而其他第三方客户端也可以直接调用此 API。

Page 28: CUDA Debugger API 的架构，展示了它如何支持各类开发和调试工具。

7. Compute Sanitizer

Compute Sanitizer 是一款用于自动扫描代码中存在的错误和内存问题的工具。它通过其子工具检查代码的正确性问题：
- Memcheck: 内存访问错误和内存泄漏检测工具。
- Racecheck: 共享内存数据访问风险检测工具。
- Initcheck: 未初始化的设备全局内存访问检测工具。
- Synccheck: 线程同步风险检测工具。

主要功能包括：
* 自动内存分配填充。
* 应用程序和内核过滤。
* Coredump 和调试器交互。
* 流序竞态条件检测。
* 未使用内存报告。

该工具还提供 API，允许开发者构建第三方开发工具。
- 示例代码库：https://github.com/NVIDIA/compute-sanitizer-samples

Page 30: Compute Sanitizer 的主要功能。

7.1 自动内存分配填充 (Automatic memory allocation padding)

该功能自 CUDA 11.3 版本引入。它通过在内存分配的末尾自动添加填充（padding），来帮助防止内存越界访问错误。

如下图所示，在没有填充的情况下，内存访问（绿色箭头）可能会越界。通过添加填充（pad区域），即使存在轻微的越界访问（红色箭头），也不会访问到不相关的内存区域，从而更容易捕获和调试此类错误。

可以通过命令行参数 --padding 32 来启用此功能，指定填充的大小。

7.2 应用和内核过滤 (Application & kernel filtering)

Compute Sanitizer 允许用户通过多种方式过滤要分析的内核或进程，以便更精确地定位问题。

按内核名称 (By kernel name):
- --kernel-regex kernel_substring=fibonacci: 仅检查名称中包含 "fibonacci" 的内核。
- --kernel-regex-exclude kernel_substring=matrixMul: 排除名称中包含 "matrixMul" 的内核。
按内核启动计数 (By kernel launch count):
- --launch-count 2 --launch-skip 2: 跳过前两次内核启动，然后分析接下来的两次内核启动。
按进程 (By process):
- --target-processes all: 分析由脚本启动的所有子进程。
- --target-processes-filter matrixMul: 仅分析名为 "matrixMul" 的子进程。

7.3 Coredump 及调试器集成

步骤 1: 生成 coredump

为了演示 coredump 的生成，我们使用一个包含内存访问错误的 CUDA 程序示例。

示例代码 (coredump_demo.cu):
该程序启动了 8 个线程，但内核 demo 中的代码 out[threadIdx.x + 2] 会导致索引为 6 和 7 的线程访问数组 out 的界外内存（访问索引 8 和 9，而数组大小为 8）。

$ cat coredump_demo.cu
static constexpr int NUM_THREADS = 8;

__global__ void demo(int *in, int *out) {
    out[threadIdx.x + 2] = in[threadIdx.x] * 5;
}

int main() {
    constexpr size_t Size = NUM_THREADS * sizeof(int);
    int* d_in = nullptr; cudaMalloc(&d_in, Size);
    int* d_out = nullptr; cudaMalloc(&d_out, Size);

    demo<<<1, NUM_THREADS>>>(d_in, d_out);
    cudaDeviceSynchronize();
}

$ nvcc -G coredump_demo.cu -o demo

使用 compute-sanitizer 运行该程序，可以检测到错误并生成一个 coredump 文件。

命令:
$ compute-sanitizer --generate-coredump yes --show-backtrace no ./demo

输出:
工具报告了两次无效的全局内存写操作，分别由线程 (6,0,0) 和 (7,0,0) 引起。同时，它生成了一个名为 core_1676502639_ubuntu_412374.nvcudmp 的 coredump 文件，并提示可以使用 cuda-gdb 加载该文件进行调试。

步骤 2: 在调试器中加载 coredump

使用 cuda-gdb 加载上一步生成的 coredump 文件，可以检查程序崩溃时的状态。

命令:
$ cuda-gdb -ex 'target cudacore core_1676502639_ubuntu_412374.nvcudmp'

调试会话:
cuda-gdb 加载 coredump 后，会自动定位到导致错误的线程和代码位置。通过 info cuda threads 命令，可以查看所有线程的状态，其中带星号（*）的行表示当前出错的线程。这使得开发者能够精确地了解 GPU 在发生错误时的执行上下文，包括线程索引、代码位置等信息。

7.4 流序竞争条件检测 (Stream-ordered race detection)

分配前使用 (Use before allocation)

在 CUDA 中，使用流（stream）可以实现异步操作。然而，不正确的同步可能导致竞争条件。下图展示了一种“分配前使用”的场景：

在 stream 1 中异步分配内存 (cudaMallocAsync)。
在 stream 2 中启动一个使用该内存的内核。
由于两个操作在不同的流上且没有同步，内核可能会在内存分配完成之前就开始执行，导致非法内存访问。
调用 cudaDeviceSynchronize() 可以确保内存分配完成后再执行后续操作，从而避免此问题。

释放后使用 (Use after free)

类似地，也可能发生“释放后使用”的竞争条件：
1. 在 stream 2 上启动一个内核，此时指针 ptr 是有效的。
2. 在 stream 1 中异步释放该内存 (cudaFreeAsync)。
3. 在 stream 2 上又启动一个使用该指针的内核。
如果第二个内核在内存被 stream 1 上的操作释放后才执行，就会导致非法内存访问。同样，需要正确的流同步来保证操作的顺序。

7.5 未使用内存报告 (Unused memory report)

compute-sanitizer 的 initcheck 工具可以报告内存分配中未被使用的部分，帮助开发者优化内存使用。

示例代码:
以下代码分配了 10 个整数的内存空间，但只启动了 8 个线程来写入数据，导致最后 2 个整数（8字节）的空间未被使用。

static constexpr int NUM_THREADS = 8;

__global__ void demo(int *in, int *out) {
    out[threadIdx.x] = in[threadIdx.x] * 5;
}

int main() {
    constexpr size_t Size = 10 * sizeof(int);
    int* d_in = nullptr; cudaMalloc(&d_in, Size);
    int* d_out = nullptr; cudaMalloc(&d_out, Size);

    demo<<<1, NUM_THREADS>>>(d_in, d_out);
    cudaDeviceSynchronize();
}

命令与输出:
使用 --track-unused-memory yes 选项运行 initcheck，工具会报告：
- 在大小为 40 字节的内存分配中，有 8 字节（从偏移量 0x20 开始）未被写入。
- 该分配中有 20% 的内存未使用。

开发者还可以通过 --unused-memory-threshold 参数设置报告的阈值。

8. 如何获取工具？

所有 NVIDIA 开发者工具都可以在官方开发者网站上找到：
https://developer.nvidia.com/tools-overview

9. GTC 相关的开发者工具资源

本页列出了在 GTC (GPU Technology Conference) 上与 NVIDIA 开发者工具相关的各种资源。

技术会议 (Sessions):
- S51205: 从宏观到微观 - CUDA 开发者工具查找并修复各种规模的问题
- S51421: 大规模优化：调查和解决多节点工作负载的隐藏瓶颈
- S51882: 使用 Source Page in Nsight Compute 更快地编写高性能 CUDA 内核
- S51772: 调试 CUDA：CUDA 正确性工具概述
- S51230: 使用 Nsight Developer Tools 进行 Orin 性能分析
- SE52434: Jetson Edge AI 开发者日：使用 NVIDIA Nsight Developer Tools 充分利用您的 Jetson Orin
实验环节 (Labs):
- DLITS1143: 使用 Nsight Profiling Tools 高效掌握通用优化模式
- DLITS1202: 调试和分析 CUDA 应用程序的正确性
- DLITS1580: 使用 NVIDIA Nsight Graphics 和 NVIDIA Nsight Systems 进行光线追踪开发
与专家交流 (Connect with Experts):
- CWES52036: 您的 CUDA 工具箱里有什么？针对最新架构的 CUDA 分析、优化和调试工具
- CWES52009: 使用 NVIDIA Developer Tools 优化光线追踪
工具下载:
- 开发者工具是免费的，并在最新的 CUDA 工具包中提供。
- https://developer.nvidia.com/cuda-downloads
支持:
- https://forums.developer.nvidia.com/c/development-tools/
更多信息:
- https://developer.nvidia.com/tools-overview