第5章：CUDA 编程、Profiling 与 Debugging

性能工程的第一性原理是 “先 profile 再优化”——这条原理在 Ch1 已经反复强调。这一章把它落到 NVIDIA 工具链上：从 CUDA 编程模型开始（理解工具看到的是什么），逐一过 Nsight Systems / Nsight Compute / cuda-gdb / compute-sanitizer / NVTX——把每个工具的”什么时候用、看什么、怎么解读”讲清楚。看完这一章你应该能拿着一个新写的 kernel 直接做完一轮性能 + 正确性的全面体检。

📑 目录

1. CUDA 编程模型快速过
2. Nsight Systems：看 CPU/GPU/通信全貌
3. Nsight Compute：单 kernel 深挖
4. cuda-gdb：调试 GPU kernel
5. compute-sanitizer：内存 / 同步错误检查
6. NVTX：给训练 step 打标签
7. 工具链的标准化使用顺序
自我检验清单
参考资料

1. CUDA 编程模型快速过

1.1 三层并行结构

CUDA 的并行抽象从粗到细有三层：

层	单位	数量级	物理对应
Grid	一次 kernel 启动	1	整次任务
Block	线程块	$10^2$ - $10^4$	调度到一个 SM 上
Thread	线程	$10^5$ - $10^9$	32 个一组成 warp

理解这三层最关键的事：Block 是”调度单位”，Warp 是”执行单位”。一个 Block 一旦被分配到 SM 上就不会被迁走（直到所有线程结束），但 Block 内部的 32 个线程组成的 warp 才是真正”同步执行”的单位。

1.2 Warp：硬件 SIMD 的真实形态

GPU 上每条指令实际上是被一个 warp 的 32 个线程同时执行的。这件事的几个直接工程后果：

branch divergence：warp 内 32 个线程如果走不同分支，硬件会串行执行所有分支——一条 if-else 让 warp 有效吞吐砍半
memory coalescing：warp 的 32 个线程同时发起访存——如果访存地址连续，硬件可以合并成一次大访存（128 字节）
__syncthreads()：Block 内同步——warp 之间的栅栏，warp 内本来就是同步的

1.3 内存层次

CUDA 内存层次比 CPU 复杂：

层	容量	延迟	谁能访问
Register	256 个 / thread	1 cycle	当前 thread
Shared Memory	48-228 KB / SM	~30 cycles	同 Block 内
L1 Cache	与 Shared 共享	~30 cycles	SM 内
L2 Cache	~50 MB（H100）	~200 cycles	全 GPU
HBM (Device Memory)	80 GB / GPU	~500 cycles	全 GPU
Host Memory	TB 级	~10K cycles	通过 PCIe / NVLink

每多一层延迟差 ~10 倍——这就是为什么 Ch6 会反复讲”shared memory tiling”的重要性。

1.4 Stream：异步并发

同一个 stream 内的操作严格按顺序
不同 stream 之间是潜在并发的——硬件会自动重叠
cudaStreamSynchronize 等待某个 stream，cudaDeviceSynchronize 等待全部

🌟 关键概念：性能工程师在 Nsight Systems 时间线上看的，主要就是每个 stream 的时间线和它们之间的重叠程度。

2. Nsight Systems：看 CPU/GPU/通信全貌

2.1 Nsight Systems 解决什么问题

Nsight Systems（简称 nsys）是全栈时间线 profiler——它把 CPU 调度、GPU kernel、CUDA API 调用、网络通信、文件 IO 全部摆在同一条时间线上。

它不告诉你”某个 kernel 内部慢在哪”——那是 Nsight Compute 的事。它告诉你”系统层面谁在等谁、谁在空闲、谁在抢资源”。

2.2 基本用法

# 最小命令：profile 一段 Python 训练脚本
nsys profile -o my_trace python train.py

# 通常加几个常用参数
nsys profile \
    --stats=true                        # 终端打印汇总
    --trace=cuda,nvtx,osrt              # 追踪范围
    --capture-range=cudaProfilerApi     # 用 API 控制起止
    -o my_trace                         # 输出文件
    python train.py

输出文件 my_trace.nsys-rep 用 Nsight Systems GUI 打开。

2.3 时间线上要看的 5 件事

打开一份 nsys 时间线后，按这个顺序看：

1. GPU 利用率（GPU 行的”色彩密度”）

全条彩色 → 利用率高
大量空白 → GPU 在等 CPU / IO / 通信

2. CPU - GPU 间的等待

CPU 调用 cudaMemcpy 后等几个 ms → 数据传输是瓶颈
CPU 调用 cudaLaunchKernel 后立即继续 → 异步 launch 正常

3. Stream 间的重叠

同时有多个 stream 跑 → 并发度好
只有一个 stream 在用 → 没充分利用 GPU 并发能力

4. NCCL / NVLink 通信带宽

通信占用比例越大，分布式扩展性越差
通信和计算是否重叠（通信发生时计算 stream 是否还在跑）

5. NVTX 标签的层次

训练 step / forward / backward / optimizer 各阶段时间分布
哪个阶段花的时间最多

2.4 一个真实例子

假设 nsys 显示这样的模式：

CPU thread 1:  [load]....[load]....[load]....
GPU stream 0:  ........[fwd] [bwd] [opt][fwd] [bwd] [opt]
GPU stream 1:                        ......

可见：

CPU 在 load → GPU 在等 → DataLoader 是瓶颈
没有 stream 1 的活动 → 没启用通信和计算并发
修复路径：增加 DataLoader workers + pin_memory + prefetch；启用 NCCL 异步通信

🌟 Nsys 的核心价值：让你直观地看见这些瓶颈，而不是靠猜。

3. Nsight Compute：单 kernel 深挖

3.1 Nsight Compute 的角色

如果说 Nsight Systems 回答”系统层面谁在等谁”，Nsight Compute（ncu）回答”某个 kernel 内部慢在哪”——它给你一份 SM 内部的硬件性能计数器报告。

3.2 基本用法

# 给某个 kernel 做完整 metrics 采集
ncu --set full -o my_kernel python train.py

# 只采几个关键 metric（速度快得多）
ncu --metrics sm__cycles_active.avg.pct_of_peak_sustained_elapsed,\
              sm__inst_executed.sum,\
              dram__bytes.sum.per_second \
    -o quick python train.py

# 只对特定 kernel 名采集
ncu --kernel-name "matmul_kernel" -o matmul python train.py

3.3 ncu 报告里的关键 metrics

打开 ncu 报告，按这个顺序看：

1. Roofline / GPU Speed Of Light

这是首页——告诉你 kernel 是 memory bound 还是 compute bound
Memory bound：DRAM bandwidth 利用率 > 80%，compute 利用率 < 30%
Compute bound：反过来
大多数深度学习 kernel 是 memory bound——优化重点是减少访存量

2. Compute Workload Analysis

SM Throughput：SM 实际利用率
Achieved Occupancy：实际 occupancy（活跃 warp / 理论最大）
低 occupancy 通常意味着 register / shared memory 用太多

3. Memory Workload Analysis

L1 Hit Rate / L2 Hit Rate
DRAM Throughput
Coalesced Access Ratio：访存合并比例（应当接近 100%）

4. Source Counters

把 metrics 投影到源代码行
看哪一行的 stall 最多（memory stall / barrier stall / pipe busy）

3.4 一个真实诊断流程

假设 ncu 显示：

DRAM Throughput 75%（接近峰值）
L2 Hit Rate 30%
这是 memory bound 模式

修复方向：

增加 shared memory tiling，让数据复用 → L2 Hit Rate 上升
减少冗余访存（合并多个相邻 kernel 共用的数据）
用 vector load（float4 / int4）减少访存指令数

ncu 报告就像 kernel 的”心电图”——看完它你就知道接下来该改哪一行代码。

4. cuda-gdb：调试 GPU kernel

4.1 cuda-gdb 解决什么问题

写 CUDA kernel 时常见的痛苦：kernel 里某个变量算错了 / 某次访存越界 / 死锁了——但 printf 在 1024 个线程里都打一遍根本看不清。cuda-gdb 是 GPU 上的 gdb——可以下断点、单步、看变量、看寄存器。

4.2 基本用法

编译时加 -g -G 让 nvcc 保留调试信息：

nvcc -g -G my_kernel.cu -o my_program
cuda-gdb ./my_program

进 cuda-gdb 后：

(cuda-gdb) break my_kernel.cu:42       # 在 my_kernel.cu 第 42 行下断点
(cuda-gdb) run                         # 运行
(cuda-gdb) cuda thread (0,0,0)          # 切到 (block 0, thread 0,0,0)
(cuda-gdb) print x                     # 查变量
(cuda-gdb) cuda block (1)              # 切到 block 1
(cuda-gdb) info cuda warps             # 看所有 warp 状态

4.3 调 GPU 死锁的常见动作

GPU 死锁经常是 __syncthreads() 在分支里被部分线程执行——cuda-gdb 检查方式：

进入死锁后按 Ctrl-C
info cuda kernels 看哪些 kernel 还在跑
info cuda threads 看每个 thread 卡在哪一行
通常会看到一部分线程卡在 __syncthreads() 一部分在 if 分支外——这就是死锁根因

4.4 一个常见坑：`-G` 让性能爆降

-G 会关闭大部分编译器优化——调试版本可能比 release 慢 10-100 倍。所以：

调试时用 -g -G
复现性能问题时用 -g（保符号但不关闭优化）
生产时不带任何调试 flag

5. compute-sanitizer：内存 / 同步错误检查

5.1 compute-sanitizer 解决什么问题

CUDA 程序最难抓的两类 bug：

内存越界：写到不该写的地方，可能延迟几千个 step 才崩
race condition：两个 warp 抢同一块 shared memory，大多数时候侥幸过

compute-sanitizer 是运行时检查器——会在每次访存 / 每个原子操作处插入检查，把这两类 bug 当场抓出来。

5.2 四种检查模式

# memcheck：内存越界 / leaked / 双重 free
compute-sanitizer --tool memcheck ./my_program

# racecheck：shared memory race
compute-sanitizer --tool racecheck ./my_program

# initcheck：未初始化访存
compute-sanitizer --tool initcheck ./my_program

# synccheck：__syncthreads 在 divergent 分支里
compute-sanitizer --tool synccheck ./my_program

5.3 典型错误形态

memcheck：

========= Invalid __global__ write of size 4 bytes
=========     at 0x90 in my_kernel
=========     by thread (15,0,0) in block (3,0,0)
=========     Address 0x... is out of bounds

racecheck：

========= Race reported between Read access at 0x... in my_kernel
=========     and Write access at 0x... in my_kernel
=========     by thread (5,0,0) and thread (12,0,0) in block (0,0,0)

报告会精确到线程 ID + block ID + 源代码行号——非常容易定位。

5.4 工程上的使用习惯

每个新 kernel 写完先跑一遍 memcheck + racecheck
CI pipeline 里加 memcheck 测试——把内存错误当回归测试做
生产前用 --tool memcheck --launch-timeout=300 跑端到端 pipeline

⭐ 关键经验：很多看似”诡异的训练发散”问题，最终被 memcheck 抓到是某个 kernel 的内存越界——浪费了几天调试，其实一个工具几分钟就能定位。

6. NVTX：给训练 step 打标签

6.1 NVTX 是什么

NVTX (NVIDIA Tools Extension) 是一组 API——让你在代码里手动插入”标签”，这些标签会出现在 Nsight Systems 时间线上。

它不影响性能（生产可保留），但极大提升 profiler 输出的可读性。

6.2 基本用法（C++）

#include <nvtx3/nvToolsExt.h>

void train_step() {
    nvtxRangePush("forward");          // 起标签
    forward(...);
    nvtxRangePop();                    // 结束

    nvtxRangePush("backward");
    backward(...);
    nvtxRangePop();

    nvtxRangePush("optimizer");
    optimize(...);
    nvtxRangePop();
}

6.3 PyTorch 集成

PyTorch 自带 torch.cuda.nvtx：

import torch.cuda.nvtx as nvtx

for step, batch in enumerate(loader):
    nvtx.range_push(f"step_{step}")
    
    nvtx.range_push("forward")
    out = model(batch)
    nvtx.range_pop()
    
    nvtx.range_push("backward")
    loss = criterion(out, target)
    loss.backward()
    nvtx.range_pop()
    
    nvtx.range_push("optimizer")
    optimizer.step()
    nvtx.range_pop()
    
    nvtx.range_pop()

6.4 标签层次设计

实际工程里推荐三层标签：

顶层（粗粒度）：epoch_X / step_Y
中层（中粒度）：forward / backward / optimizer / data_loading
底层（细粒度）：layer_attention / layer_mlp / kernel_softmax 等

层次清晰的 NVTX 会让 nsys 时间线像目录树一样可折叠，定位问题事半功倍。

7. 工具链的标准化使用顺序

把上面四个工具串成一个标准 SOP——拿到一个新 kernel / 新模型时按这个顺序走：

7.1 Step 1：用 NVTX 打标签

让训练 / 推理代码每个关键阶段都被 NVTX 包起来。这是永久性投入——生产代码里也保留。

7.2 Step 2：用 nsys 看时间线

nsys profile --stats=true -o overview python train.py

回答：

GPU 利用率多少？
是 CPU 等 GPU 还是 GPU 等 CPU？
通信和计算重叠吗？
哪个阶段最慢？

如果系统层面瓶颈明显（比如 DataLoader 慢、通信慢），先解决系统层面——这一步不需要 ncu。

7.3 Step 3：用 ncu 深挖最慢的 kernel

nsys 报告里找到耗时最多的几个 kernel，对它们做：

ncu --set full --kernel-name "regex" -o detail python train.py

回答：

Memory bound 还是 Compute bound？
Occupancy 多高？
L1/L2 Hit Rate 多高？
Source Counter 显示哪一行 stall 最多？

7.4 Step 4：写新 kernel 前先 sanity 检查

新 kernel 写完，先：

compute-sanitizer --tool memcheck ./test
compute-sanitizer --tool racecheck ./test

通过后再做性能优化——避免在错误代码上做”假优化”。

7.5 Step 5：debug 用 cuda-gdb

发现 numerical 错误 / 死锁等具体 bug 时再用 cuda-gdb——多数性能问题不需要它。

7.6 这套 SOP 的工程价值

很多团队的常见误区是**“先优化再 profile”**——凭直觉改代码，结果优化错地方。这套 SOP 强制要求：

✅ 先看大图（nsys）
✅ 再看局部（ncu）
✅ 改代码前先确保正确性（sanitizer）
✅ 永远基于真实数字决策

按这个 SOP，性能优化能从艺术变成工程。

🎯 自我检验清单

CUDA 三层并行（Grid/Block/Thread）和硬件（GPU/SM/Warp）的对应关系是什么？为什么 Block 是”调度单位”而 Warp 是”执行单位”？
Branch divergence 在硬件上怎么发生？为什么它对性能影响大？
Nsight Systems 和 Nsight Compute 的角色分别是什么？什么时候该用哪个？
compute-sanitizer 的 4 种检查模式各自检查什么？为什么 racecheck 在普通测试中难以触发？
Nvtx 标签的三层设计在 nsys 时间线上各自服务什么 debug 目的？

📚 参考资料

CUDA C++ Programming Guide：docs.nvidia.com/cuda/cuda-c-programming-guide/
NVIDIA Nsight Systems：docs.nvidia.com/nsight-systems/
NVIDIA Nsight Compute：docs.nvidia.com/nsight-compute/
cuda-gdb 文档：docs.nvidia.com/cuda/cuda-gdb/
compute-sanitizer 文档：docs.nvidia.com/compute-sanitizer/
NVTX 文档：github.com/NVIDIA/NVTX
本系列模块二：CUDA 编程与算子优化（详细 kernel 写作教程）

下一章预告：Ch6 把这一章建立的工具链能力用在单 kernel 深度调优——memory coalescing / occupancy / shared memory tiling / TMA / WGMMA / kernel fusion 一条线串起来，给出”朴素 GEMM → cuBLAS 级”7 步演进的实战路径。

搜索