cuda

CUDA 编程技能

核心理念

先测量，再猜测。 GPU 性能往往很反直觉。先分析，再假设，然后修改，最后验证。

改动要小且相互独立。 CUDA 错误会叠加。一次只做一个改动，测试后再提交。不要急着“一次性修复所有问题”。

printf 是你最有力的工具。 调试器失效、工具输出难以理解时，设备代码中的 printf 能揭示真相。大量使用也不必觉得不好意思。

有时，仔细查看差异就够了。 难以理解的段错误很常见，工具往往也帮不上忙。人工处理方法：尽量缩小差异，仔细阅读，找出错误。这完全合理，而且通常比使用工具更快。

调试工作流

遇到错误时的第一步

1. 尽可能精简地复现 — 使用尽可能小的输入隔离出故障内核
2. 添加 printf — 使用任何工具前，先在设备代码中添加 printf 来跟踪执行过程
3. 运行 compute-sanitizer — 以非交互方式捕获内存错误：

   compute-sanitizer --tool memcheck ./your_program
   compute-sanitizer --tool racecheck ./your_program  # for race conditions
   compute-sanitizer --tool initcheck ./your_program  # uninitialized memory
   

4. 如果仍然卡住，尝试以非交互方式使用 cuda-gdb 获取回溯：

   cuda-gdb -batch -ex "run" -ex "bt" ./your_program
   

5. 工具失效时 — 尽量缩小正常代码与故障代码之间的差异。仔细阅读。错误就在差异中。

设备代码中的 printf

__global__ void myKernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx == 0) {  // Limit output
        printf("Kernel launched, n=%d, data[0]=%f\n", n, data[0]);
    }
    // ... kernel logic ...
    if (idx < 10) {  // Sample a few threads
        printf("Thread %d: result=%f\n", idx, someValue);
    }
}

关键模式：

• 使用 if (idx == 0) 或 if (idx < N) 设置条件，避免输出过多
• 在内核入口处打印，以确认内核已启动
• 在怀疑出错的位置打印中间值
• 内核完成时会自动刷新

compute-sanitizer 快速参考

常见陷阱：“Invalid shared write... out of bounds” 通常意味着内核启动时动态共享内存分配不足，而不是数组索引错误。请检查 <<<grid, block, smem_size>>>。

# Memory errors (most common)
compute-sanitizer --tool memcheck ./program

# Other tools: racecheck, initcheck, synccheck
# For detailed options, see references/debugging-tools.md

非交互式 cuda-gdb

# Get backtrace on crash
cuda-gdb -batch -ex "run" -ex "bt" ./program

# For breakpoints, thread inspection, see references/debugging-tools.md

使用调试信息编译：

nvcc -g -G -lineinfo program.cu -o program

使用 cuobjdump 检查二进制文件

# Dump PTX and SASS
cuobjdump -ptx ./program
cuobjdump -sass ./program

# For resource usage, symbol listing, see references/debugging-tools.md

完整的调试工具参考： 详细了解 compute-sanitizer 选项、cuda-gdb 工作流和 cuobjdump 分析模式，请参阅 references/debugging-tools.md。

性能优化工作流

黄金法则

未经性能分析，绝不进行优化。 对 GPU 瓶颈的直觉几乎总是错的。“分析 → 修复 → 验证”循环才是真正的优化工作，而非准备步骤。

性能排查步骤

1. 建立基线 — 测量操作耗时并记录
2. 使用 nsys 分析 — 获取时间线，找出关键内核
3. 使用 ncu 深入分析 — 分析具体的瓶颈内核
4. 提出假设 — 根据指标提出明确的假设
5. 一次只改一处 — 进行单项、有针对性的修改
6. 验证 — 重新分析，确认性能有所提升
7. 重复

nsys（Nsight Systems）— 时间线分析

使用 nsys 回答“时间都花在哪里了？”——分析 CPU/GPU 交互、内核启动模式、内存传输和整体时间线。

# Basic profile
nsys profile -o report ./program
nsys stats report.nsys-rep --report cuda_gpu_kern_sum

# With NVTX markers
nsys profile --trace=cuda,nvtx -o report ./program

# Key reports: cuda_gpu_kern_sum, cuda_api_sum, cuda_gpu_mem_time_sum, nvtx_sum
# For detailed usage, see references/nsys-guide.md

nsys 分析模式详解： 关于时间线解读、常见瓶颈识别和分析工作流，请参阅 references/nsys-guide.md。

ncu（Nsight Compute）— 内核分析

使用 ncu 回答“这个内核为什么这么慢？”——分析详细指标、Roofline、内存和占用率。

# Profile specific kernel
ncu --kernel-name "myKernel" -o report ./program

# Quick summary to stdout
ncu --set basic ./program

# Sets: basic, full, memory, launch, roofline
# Sections: ComputeWorkloadAnalysis, MemoryWorkloadAnalysis, Occupancy
# For detailed metrics and interpretation, see references/ncu-guide.md

警告： ncu 专家系统的建议可能会误导。务必结合实际指标和实验进行验证。

规模很重要： 对大规模场景有帮助的优化可能会降低小规模场景的性能。务必在实际问题规模下进行分析，而非理论最大规模。

ncu 指标解读详解： 关于理解 Roofline 分析、内存瓶颈、占用率限制和 Warp 调度，请参阅 references/ncu-guide.md。

使用 NVTX 进行自定义插桩

需要比内核级别更细粒度的分析时，请使用 NVTX：

#include <nvtx3/nvToolsExt.h>

nvtxRangePush("Operation Name");
// ... code to profile ...
nvtxRangePop();

编译： -lnvToolsExt | 分析： nsys profile --trace=cuda,nvtx

完整模式： 关于嵌套范围、颜色和分析工作流，请参阅 references/nvtx-patterns.md。

常见性能模式

症状	可能原因	调查方法
GPU 利用率低	内核启动开销、CPU 瓶颈	查看 nsys 时间线，寻找间隙
受内存限制	访问模式不佳、缓存命中率低	查看 ncu 内存部分，检查合并访问
受计算限制但速度慢	占用率低、寄存器压力大	查看 ncu 占用率，减少寄存器使用
大量小内核	启动开销占主导	查看 nsys 时间线，考虑融合
memcpy 时间长	H2D/D2H 传输过多	查看 nsys cuda_gpu_mem，批量传输
大部分周期处于停滞状态	存储体冲突、内存停滞	查看 ncu SchedulerStatistics，检查共享内存
每次请求的扇区数偏高	合并访问不佳（每次请求 >4 个扇区）	查看 ncu 内存指标，使用向量化加载

关键陷阱： 存储体冲突和内存合并访问问题往往是性能瓶颈，但不经分析很难发现。详细诊断和修复方法请参阅 references/performance-traps.md。

现实检验： 将 80% 的优化时间预留给你未预料到的问题。基于性能分析的迭代有助于发现真正的瓶颈。

编译参考

# Debug build
nvcc -g -G -lineinfo -O0 program.cu -o program_debug

# Release build
nvcc -O3 -lineinfo program.cu -o program

# Specific architecture
nvcc -arch=sm_80 program.cu -o program  # Ampere
nvcc -arch=sm_89 program.cu -o program  # Ada Lovelace
nvcc -arch=sm_90 program.cu -o program  # Hopper

# Generate PTX (inspect it)
nvcc -ptx program.cu

# Verbose compilation (see register usage)
nvcc --ptxas-options=-v program.cu

# With NVTX
nvcc program.cu -lnvToolsExt -o program

生产环境性能分析时，始终使用 -lineinfo 编译 — 开销极低，还能关联到源代码。

本地 API 文档

可通过 grep 搜索完整参考文档：

PTX ISA 9.1 — references/ptx-docs/（405 个文件，2.3MB）

• 搜索指南：references/ptx-isa.md
• 用途：指令级优化、内联 PTX、TensorCore 操作（WMMA、WGMMA、TMA）、内存交错

CUDA Runtime API 13.1 — references/cuda-runtime-docs/（107 个文件，0.9MB）

• 搜索指南：references/cuda-runtime.md
• 用途：错误代码、API 参数、设备属性（cudaDeviceProp）、内存管理、流行为

CUDA Driver API 13.1 — references/cuda-driver-docs/（128 个文件，0.8MB）

• 搜索指南：references/cuda-driver.md
• 用途：上下文管理（cuCtxCreate）、模块加载（cuModuleLoad）、虚拟内存、Driver 错误（CUDA_ERROR_*）、高级功能

每份搜索指南都包含 grep 示例、文档结构和常见用法模式。

搜索策略： 使用 grep/ripgrep 直接搜索 *-docs/ 目录。搜索指南（.md 文件）提供导航模式和常见查询。

其他参考资料

• references/performance-traps.md — 存储体冲突、内存合并访问、依赖规模的优化
• references/debugging-tools.md — compute-sanitizer、cuda-gdb、cuobjdump 的详细用法
• references/nsys-guide.md — nsys 时间线分析和瓶颈识别
• references/ncu-guide.md — ncu 指标、屋顶线、占用率解读
• references/nvtx-patterns.md — NVTX 插桩和性能分析模式

优化前检查清单

• 已建立可复现的基准计时
• 已使用 nsys 进行性能分析并识别热点
• 知道哪些内核主导运行时间
• 已使用 ncu 对目标内核进行性能分析
• 已识别具体瓶颈（内存？计算？延迟？）
• 已提出具体、可测试的假设
• 计划只改动一项内容