Debugging CUDA: An Overview of CUDA Correctness Tools

Steve Ulrich & Aurelien Chartier | GTC March 2023

目录

1. 调试器工具

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

Page 2
Page 2

1.1 调试器工具概览

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

Page 3
Page 3

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

Page 4
Page 4

2. CUDA GDB

2.1 概述 (Overview)

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

Page 5
Page 5

2.2 架构支持 (Architecture Support)

Page 6
Page 6

2.3 基础操作 (Basics)

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

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

$ cuda-gdb --quiet
(cuda-gdb) attach 261230
Page 7
Page 7
Page 8 a `cuda-gdb` command output showing the `help info cuda` command.
Page 8 a `cuda-gdb` command output showing the `help info cuda` command.
(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]
Page 9
Page 9
(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 反汇编。

Page 10
Page 10

2.4 Coredumps

$ ./memexceptions 1
SM version: 86, Min version: 35, Max version: 999
Aborted (core dumped)
$ ls | grep core
core_1669651659_agontarek-dt_612954.nvcudmp
$ export CUDA_COREDUMP_FILE="/lus/grand/projects/alcf_training/$USER/core.gpu.%h.%p"
Page 11
Page 11

2.5 多 GPU 调试 (Multi-GPU Debugging)

Page 12
Page 12

2.6 Python 支持 (Python Support)

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

2.7 WSL2 支持 (WSL2 Support)

Page 14
Page 14

3. Nsight Visual Studio Edition

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

Page 15
Page 15

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

Page 16: Nsight Visual Studio Edition 界面展示了GPU寄存器、本地变量、断点和CUDA C++源代码。
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 或两者。
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。
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 中所有线程的详细信息。
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 架构支持。
Page 20: Nsight Visual Studio Edition 的 GPU 架构支持。

4. Nsight Visual Studio Code Edition

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

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

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

Page 22: Nsight Visual Studio Code Edition 调试界面。
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 的安装选项。
Page 23: Nsight Visual Studio Code Edition 的安装选项。

5. Nsight Eclipse Edition

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

Page 24: Nsight Eclipse Edition 的特性。
Page 24: Nsight Eclipse Edition 的特性。

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

Page 25: Nsight Eclipse Edition 的调试界面。
Page 25: Nsight Eclipse Edition 的调试界面。

5.1 安装与配置

Page 26: 在 Eclipse 中安装 Nsight 插件的步骤。
Page 26: 在 Eclipse 中安装 Nsight 插件的步骤。

5.2 Java 支持

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

6. CUDA Debugger API

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

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

Page 28: CUDA Debugger 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 的主要功能。
Page 31

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

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

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

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

Page 32
Page 32

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

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

Page 33
Page 33

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
Page 34
Page 34

使用 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 加载该文件进行调试。

Page 35
Page 35

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

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

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

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

Page 36
Page 36

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

分配前使用 (Use before allocation)

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

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

释放后使用 (Use after free)

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

Page 38
Page 38

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 参数设置报告的阈值。

Page 39
Page 39

8. 如何获取工具？

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

Page 29: NVIDIA 开发者工具概览页面。
Page 29: NVIDIA 开发者工具概览页面。

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

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