cuda-profiling

CUDA 性能分析

目的

指导 Agent 使用 Nsight Systems（时间线级）和 Nsight Compute（内核级指标）分析 CUDA 应用，包括使用 NCU CLI 自动收集指标、解读 Roofline 模型，以及诊断内核是受内存还是计算限制。

适用场景

• CUDA 内核运行速度低于预期，需要定位瓶颈
• 比较不同内核变体（分块策略、块大小）
• 使用 ncu 指标构建 CI 性能回归检查
• 关联多流流水线中的 CPU 和 GPU 活动
• 使用 NVTX 标注应用阶段，以便在时间线上查看
• 解读占用率、内存吞吐量和 SM 利用率指标

工作流

1. 选择性能分析工具

What do you need?
├── System-wide timeline (CPU+GPU+CUDA API) → Nsight Systems (nsys)
├── Per-kernel deep metrics (occupancy, memory) → Nsight Compute (ncu)
└── Quick metric from CLI in CI → ncu --metrics ...

2. Nsight Systems — 时间线性能分析

# Profile entire application
nsys profile --trace=cuda,nvtx,osrt --output=report ./my_cuda_app

# Open report
nsys-ui report.nsys-rep

# CLI summary
nsys stats report.nsys-rep

时间线上需要关注的内容：

• 内核启动之间的间隙（CPU 瓶颈或同步点）
• cudaDeviceSynchronize 导致的停顿
• 跨流的 H2D 复制与内核执行之间的重叠
• CUDA API 调用开销

# Capture with CUDA graph info
nsys profile --capture-range=cudaProfilerApi ./my_cuda_app

3. NVTX 范围注释

#include <nvtx3/nvToolsExt.h>

void pipeline(void) {
    nvtxRangePushA("H2D copy");
    cudaMemcpyAsync(d_in, h_in, size, cudaMemcpyHostToDevice, stream);
    nvtxRangePop();

    nvtxRangePushA("kernel");
    my_kernel<<<grid, block, 0, stream>>>(d_in, d_out, n);
    nvtxRangePop();

    nvtxRangePushA("D2H copy");
    cudaMemcpyAsync(h_out, d_out, size, cudaMemcpyDeviceToHost, stream);
    nvtxRangePop();
}

使用 -lnvToolsExt 编译，或链接仅包含头文件的 nvtx3。范围会在 Nsight Systems 中显示为彩色带。

4. Nsight Compute — 内核分析

# Profile all kernels, save report
ncu -o kernel_report ./my_cuda_app

# Profile specific kernel by name
ncu --kernel-name regex:matmul_tiled ./my_cuda_app

# Launch UI
ncu-ui kernel_report.ncu-rep

NCU 报告中的关键部分：

• Speed of Light：SM 吞吐量、内存吞吐量与峰值的比较
• Occupancy：活动 warp 数与硬件上限的比较
• Memory Workload Analysis：L1/L2 命中率、合并访问效率
• Warp State Statistics：停顿原因（内存、屏障、调度）

5. NCU CLI 指标

# Essential metrics set
ncu --metrics \
  sm__throughput.avg.pct_of_peak_sustained_elapsed,\
  dram__throughput.avg.pct_of_peak_sustained_elapsed,\
  sm__warps_active.avg.pct_of_peak_sustained_active,\
  l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\
  smsp__sass_thread_inst_executed_op_ffma_pred_on.sum \
  ./my_cuda_app

# CSV export for CI
ncu --csv --metrics dram__bytes_read.sum,dram__bytes_write.sum ./my_cuda_app

# Set kernel replay mode for accurate counters
ncu --kernel-replay-mode application ./my_cuda_app

6. 诊断内存受限还是计算受限

Roofline interpretation
├── dram__throughput near peak AND sm__throughput low → memory-bound
│   └── Fix: coalescing, shared mem tiling, reduce traffic
├── sm__throughput near peak AND dram low → compute-bound
│   └── Fix: tensor cores, loop unrolling, ILP
└── Both low → launch config, occupancy, or sync overhead

Roofline 模型（概念示意）：

Performance (GFLOP/s)
    |     /\  compute roof
    |    /  \
    |   /    \____ memory roof (bandwidth-limited region)
    |  /
    +------------------ Arithmetic Intensity (FLOP/byte)

测量算术强度：smsp__sass_thread_inst_executed_op_ffma_pred_on.sum * 2 / dram__bytes.sum

7. 占用率分析

ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active,\
launch__occupancy_limit_registers,\
launch__occupancy_limit_shared_mem,\
launch__occupancy_limit_block_size \
./my_cuda_app

限制因素	常见修复方法
寄存器	-maxrregcount，简化内核
共享内存	减小分块大小，拆分阶段
块大小	尝试使用 128 或 256，而不是 512+

8. 性能分析工作流检查清单

# 1. Build with line info (not -G unless debugging)
nvcc -lineinfo -O3 -arch=sm_80 -o app main.cu

# 2. Timeline first
nsys profile --trace=cuda,nvtx -o timeline ./app

# 3. Deep dive on hot kernel
ncu --kernel-name regex:hot_kernel --set full ./app

# 4. Compare before/after
ncu --csv --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed ./app_v1 > v1.csv
ncu --csv --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed ./app_v2 > v2.csv

常见问题

症状	原因	修复方法
ERR_NVGPUCTRPERM	性能分析权限不足	使用 sudo 运行，或设置 NVreg_RestrictProfilingToAdminUsers=0
所有指标均为零	性能分析已禁用或 GPU 不正确	检查 CUDA_VISIBLE_DEVICES；使用 --target-processes all
NCU 报告为空	内核执行时间过短或未启动	增加工作负载；检查 cudaGetLastError()
性能分析开销过大	对许多内核收集了完整指标集	使用 --kernel-name 筛选；使用 --launch-skip
时间线上看不到重叠	使用了单一默认流	创建多个流；使用异步复制
占用率看起来正常，但内核运行缓慢	内存延迟未被隐藏	检查内存合并访问；增加活动 warp 数

相关技能

• skills/gpu/cuda — 内核编写、占用率调优、nvcc 标志
• skills/gpu/gpu-memory-model — 合并访问、bank 冲突、SIMT 模型
• skills/gpu/cuda-debugging — 在性能调优前确保正确性
• skills/profilers/intel-vtune-amd-uprof — CPU 侧 Roofline 和热点分析
• skills/profilers/flamegraphs — 与 nsys 时间线互补的 CPU 火焰图
• skills/profilers/hardware-counters — 通用 perf stat 概念