Back to Details

cuda-profiling

Compare original and translation side by side

🇺🇸

Original

English
🇨🇳

Translation

Chinese

CUDA Profiling

CUDA Profiling

Purpose

用途

Guide agents through profiling CUDA applications with Nsight Systems (timeline-level) and Nsight Compute (kernel-level metrics), using the NCU CLI for automated metric collection, interpreting roofline models, and diagnosing whether kernels are memory-bound or compute-bound.
指导Agent使用Nsight Systems（时间线层面）和Nsight Compute（kernel级指标）对CUDA应用进行profiling，使用NCU CLI实现自动化指标采集，解读roofline模型，并诊断kernel是内存受限还是计算受限。

When to Use

适用场景

  • A CUDA kernel is slower than expected and you need bottleneck identification
  • Comparing kernel variants (tiling strategies, block sizes)
  • Building CI performance regression checks with 
    ncu
    metrics
  • Correlating CPU and GPU activity in multi-stream pipelines
  • Annotating application phases with NVTX for timeline visibility
  • Interpreting occupancy, memory throughput, and SM utilization metrics
  • CUDA kernel运行速度低于预期，需要识别瓶颈
  • 对比不同的kernel实现方案（分块策略、block大小）
  • 使用
    ncu
    指标构建CI性能回归检查
  • 在多流流水线中关联CPU和GPU的活动
  • 使用NVTX标注应用阶段，提升时间线可读性
  • 解读占用率、内存吞吐量和SM利用率指标

Workflow

工作流

1. Choose profiling tool

1. 选择profiling工具

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 ...
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 — timeline profiling

2. Nsight Systems — 时间线profiling

bash
undefined
bash
undefined

Profile entire application

Profile entire application

nsys profile --trace=cuda,nvtx,osrt --output=report ./my_cuda_app
nsys profile --trace=cuda,nvtx,osrt --output=report ./my_cuda_app

Open report

Open report

nsys-ui report.nsys-rep
nsys-ui report.nsys-rep

CLI summary

CLI summary

nsys stats report.nsys-rep

What to look for in the timeline:
- Gaps between kernel launches (CPU bottleneck or sync points)
- `cudaDeviceSynchronize` stalls
- Overlap between H2D copies and kernel execution across streams
- CUDA API call overhead

```bash
nsys stats report.nsys-rep

在时间线中需要关注的点：
- kernel启动之间的间隙（CPU瓶颈或同步点）
- `cudaDeviceSynchronize` 停顿
- 跨流的H2D拷贝与kernel执行之间的重叠
- CUDA API调用开销

```bash

Capture with CUDA graph info

Capture with CUDA graph info

nsys profile --capture-range=cudaProfilerApi ./my_cuda_app
undefined
nsys profile --capture-range=cudaProfilerApi ./my_cuda_app
undefined

3. NVTX range annotations

3. NVTX范围标注

cpp
#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();
}
Compile with 
-lnvToolsExt
or link nvtx3 header-only. Ranges appear as colored bands in Nsight Systems.
cpp
#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 — kernel analysis

4. Nsight Compute — kernel分析

bash
undefined
bash
undefined

Profile all kernels, save report

Profile all kernels, save report

ncu -o kernel_report ./my_cuda_app
ncu -o kernel_report ./my_cuda_app

Profile specific kernel by name

Profile specific kernel by name

ncu --kernel-name regex:matmul_tiled ./my_cuda_app
ncu --kernel-name regex:matmul_tiled ./my_cuda_app

Launch UI

Launch UI

ncu-ui kernel_report.ncu-rep

Key sections in NCU report:
- **Speed of Light**: SM throughput vs memory throughput vs peak
- **Occupancy**: Active warps vs hardware limit
- **Memory Workload Analysis**: L1/L2 hit rates, coalescing efficiency
- **Warp State Statistics**: Stall reasons (memory, barrier, dispatch)
ncu-ui kernel_report.ncu-rep

NCU报告中的关键部分：
- **Speed of Light**：SM吞吐量、内存吞吐量与峰值的对比
- **Occupancy**：活跃warp与硬件上限的对比
- **内存工作负载分析**：L1/L2命中率、合并效率
- **Warp状态统计**：停顿原因（内存、屏障、调度）

5. NCU CLI metrics

5. NCU CLI指标

bash
undefined
bash
undefined

Essential metrics set

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

CSV export for CI

ncu --csv --metrics dram__bytes_read.sum,dram__bytes_write.sum ./my_cuda_app
ncu --csv --metrics dram__bytes_read.sum,dram__bytes_write.sum ./my_cuda_app

Set kernel replay mode for accurate counters

Set kernel replay mode for accurate counters

ncu --kernel-replay-mode application ./my_cuda_app
undefined
ncu --kernel-replay-mode application ./my_cuda_app
undefined

6. Memory-bound vs compute-bound diagnosis

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 model (conceptual):
Performance (GFLOP/s)
    |     /\  compute roof
    |    /  \
    |   /    \____ memory roof (bandwidth-limited region)
    |  /
    +------------------ Arithmetic Intensity (FLOP/byte)
Measure arithmetic intensity: 
smsp__sass_thread_inst_executed_op_ffma_pred_on.sum * 2 / dram__bytes.sum
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. Occupancy analysis

7. 占用率分析

bash
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
Limiting factorTypical fix
Registers
-maxrregcount
, simplify kernel
Shared memoryReduce tile size, split phases
Block sizeTry 128 or 256 instead of 512+
bash
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
限制因素典型优化方案
寄存器（Registers）
-maxrregcount
、简化kernel
共享内存（Shared memory）减小分块大小、拆分阶段
Block大小尝试128或256，而非512及以上

8. Profiling workflow checklist

8. Profiling工作流检查清单

bash
undefined
bash
undefined

1. Build with line info (not -G unless debugging)

1. Build with line info (not -G unless debugging)

nvcc -lineinfo -O3 -arch=sm_80 -o app main.cu
nvcc -lineinfo -O3 -arch=sm_80 -o app main.cu

2. Timeline first

2. Timeline first

nsys profile --trace=cuda,nvtx -o timeline ./app
nsys profile --trace=cuda,nvtx -o timeline ./app

3. Deep dive on hot kernel

3. Deep dive on hot kernel

ncu --kernel-name regex:hot_kernel --set full ./app
ncu --kernel-name regex:hot_kernel --set full ./app

4. Compare before/after

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
undefined
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
undefined

Common Problems

常见问题

SymptomCauseFix
ERR_NVGPUCTRPERM
Insufficient profiling permissionsRun with sudo or set 
NVreg_RestrictProfilingToAdminUsers=0
All metrics show zeroProfiling disabled or wrong GPUCheck 
CUDA_VISIBLE_DEVICES
; use 
--target-processes all
NCU report emptyKernel too short or not launchedIncrease workload; verify 
cudaGetLastError()
Huge profiling overheadFull metric sets on many kernelsUse 
--kernel-name
filter; 
--launch-skip
Timeline shows no overlapSingle default streamCreate multiple streams; use async copies
Occupancy looks fine but kernel slowMemory latency not hiddenCheck memory coalescing; increase active warps
现象原因解决方案
ERR_NVGPUCTRPERM
profiling权限不足使用sudo运行，或设置
NVreg_RestrictProfilingToAdminUsers=0
所有指标显示为零profiling被禁用或选择了错误的GPU检查
CUDA_VISIBLE_DEVICES
；使用
--target-processes all
参数
NCU报告为空Kernel运行时间过短或未启动增加工作负载；验证
cudaGetLastError()
返回值
Profiling开销过大对大量kernel采集全量指标使用
--kernel-name
过滤；添加
--launch-skip
参数
时间线无重叠使用单个默认流创建多个流；使用异步拷贝
占用率正常但kernel运行缓慢内存延迟未被隐藏检查内存合并；增加活跃warp数量

Related Skills

相关技能

  • skills/gpu/cuda
    — kernel writing, occupancy tuning, nvcc flags
  • skills/gpu/gpu-memory-model
    — coalescing, bank conflicts, SIMT model
  • skills/gpu/cuda-debugging
    — correctness before performance tuning
  • skills/profilers/intel-vtune-amd-uprof
    — CPU-side roofline and hotspot analysis
  • skills/profilers/flamegraphs
    — CPU flamegraphs complementary to nsys timeline
  • skills/profilers/hardware-counters
    — general perf stat concepts
  • skills/gpu/cuda
    — kernel编写、占用率调优、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
    — 通用性能计数器概念