📚 课程目录第4章开发环境搭建与调试工具

4. 开发环境搭建与调试工具

CUDA 自定义算子的开发效率，很大程度上取决于工具链是否完整、调试手段是否到位。一个可用的环境至少包含四层：驱动与运行时、编译器 nvcc、系统级性能分析工具（Nsight Systems）、内核级性能分析工具（Nsight Compute），以及调试与正确性检查工具（cuda-gdb、compute-sanitizer）。本节按“装—编—测—调—析”的顺序，把这条链路讲清楚，并给出可直接执行的命令与判读方法。

一、CUDA Toolkit 安装与版本对齐

安装前必须先确认三件事：GPU 架构、驱动版本、目标 CUDA 版本。驱动决定了可支持的最高 CUDA 运行时版本，例如驱动 535.x 大致对应 CUDA 12.2 上限；若 Toolkit 版本高于驱动支持范围，会出现 cudaErrorInsufficientDriver。查看方式：

nvidia-smi                      # 查看驱动版本、CUDA Version 上限、GPU 型号
nvcc --version                  # 查看已安装 Toolkit 版本
ls /usr/local | grep cuda       # 查看本机已安装的多个 Toolkit

推荐使用官方 runfile 或 apt 仓库安装，避免与系统自带驱动冲突。以 Ubuntu + CUDA 12.4 为例：

wget https://developer.download.nvidia.com/compute/cuda/12.4.0/local_installers/cuda_12.4.0_550.54.14_linux.run
sudo sh cuda_12.4.0_550.54.14_linux.run --toolkit --silent --override
echo 'export PATH=/usr/local/cuda-12.4/bin:$PATH' >> ~/.bashrc
echo 'export LD_LIBRARY_PATH=/usr/local/cuda-12.4/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc
source ~/.bashrc

安装后务必验证：nvcc --version 输出应显示 release 12.4；nvidia-smi 顶部 CUDA Version 应不低于 12.4。若机器上存在多个 Toolkit，用 update-alternatives 或直接切换 PATH 中的软链接目录来指定默认版本，避免编译与运行库版本错配。

二、nvcc 编译：从源码到可执行内核

nvcc 是 CUDA 的编译驱动，它把 .cu 文件拆分为主机代码与设备代码，分别交给 gcc 和 cicc/ptxas 处理，最后链接成可执行文件。理解它的关键参数，是后续性能分析的前提。

nvcc -O3 -arch=sm_80 -lineinfo -Xptxas -v \
     -o vector_add vector_add.cu

# -arch=sm_80       目标架构（A100），决定生成 SASS 的指令集
# -lineinfo         保留源码行号，Nsight Compute 定位热点必需
# -Xptxas -v        输出寄存器、共享内存、spill 统计
# -G                生成设备调试信息（配合 cuda-gdb，会禁用优化）

编译后 -Xptxas -v 会打印类似 Used 24 registers, 0 bytes spill stores 的信息。寄存器数直接决定 occupancy：若某内核每线程用 64 个寄存器，在 64K 寄存器/SM 的架构上，单 SM 最多容纳 1024 线程，occupancy 上限约 50%。出现 spill 说明寄存器压力过大，通常需要减小 unroll 或拆分内核。

三、调试与正确性检查

内核崩溃、结果错误、越界访问是三类最常见问题，对应不同工具。

cuda-gdb 用于定位崩溃点与逻辑错误。编译时加 -g -G，运行：

cuda-gdb ./vector_add
(cuda-gdb) set cuda memcheck on
(cuda-gdb) run
(cuda-gdb) cuda thread              # 查看当前线程坐标
(cuda-gdb) print idx                # 打印设备变量
(cuda-gdb) cuda block 0 thread 5    # 切换到指定线程

compute-sanitizer 是 memcheck 的替代品，用于检测越界、未初始化读取、竞态。它不需要重新编译，直接包裹可执行文件：

compute-sanitizer --tool memcheck ./vector_add
compute-sanitizer --tool racecheck ./vector_add   # 共享内存竞态
compute-sanitizer --tool initcheck ./vector_add   # 未初始化访问

典型输出会指出 Invalid __global__ read of size 4 并给出地址与线程坐标，据此可直接定位到数组下标越界的行。

四、性能分析：Nsight Systems 与 Nsight Compute

两者分工明确：Nsight Systems 看“时间花在哪”，Nsight Compute 看“内核为什么慢”。

nsys profile -o report ./vector_add
nsys stats report.nsys-rep

ncu --set full -o kernel_report ./vector_add
ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed \
    --metrics gpu__time_duration.sum ./vector_add

Nsight Systems 的时间线会显示 H2D 拷贝、内核执行、D2H 拷贝的占比。若拷贝占 70%，说明瓶颈在 PCIe 传输，应考虑 pinned memory 或减少传输次数。Nsight Compute 则给出每个内核的 SM 利用率、内存吞吐、warp 停顿原因。判读时优先看三个指标：sm__throughput（计算单元是否吃满）、dram__throughput（显存带宽是否吃满）、smsp__warp_issue_stalled_*（warp 为何停顿）。若计算与显存吞吐都低，通常是 latency bound，需要提高 occupancy 或增加 ILP。

五、性能指标解读要点

拿到 Nsight Compute 报告后，不要被几十个指标淹没，按以下顺序判读即可：

指标	含义	判读与动作
sm__throughput.avg.pct_of_peak_sustained_elapsed	SM 计算单元利用率	>80% 说明计算密集，考虑算法优化；低则看内存
dram__throughput.avg.pct_of_peak_sustained_elapsed	显存带宽利用率	>80% 说明带宽受限，考虑合并访问、减少冗余读取
smsp__warp_issue_stalled_long_scoreboard	等待全局内存返回	高则增加 ILP 或预取，提高 occupancy
smsp__warp_issue_stalled_barrier	等待 __syncthreads	高则检查负载均衡与同步次数
launch__occupancy_limit_registers	寄存器限制的 occupancy	若明显低于理论值，用 __launch_bounds__ 限制寄存器

建议把 -lineinfo 作为日常编译的默认选项。它只增加少量编译时间，却能让 Nsight Compute 把指标精确映射到源码行，否则你只能看到 SASS 地址，定位成本会成倍上升。

使用 -G 编译会禁用大部分优化，生成的 SASS 与发布版本差异巨大，性能数据不可用于最终结论。正确做法是：调试用 -g -G，性能分析用 -O3 -lineinfo，两套构建分开。

把上述流程串起来，一个典型的内核开发闭环是：先用 nvcc 编译并查看寄存器/spill 报告，再用 compute-sanitizer 确认无越界与竞态，用 cuda-gdb 修掉逻辑错误，最后用 Nsight Systems 定位系统级瓶颈、用 Nsight Compute 定位内核级瓶颈，依据指标决定是改访存模式、调 occupancy 还是改算法。工具本身不产生性能，但它们把“猜”变成“看”，这是自定义算子优化能否持续迭代的关键。