使用方法
- 复制安装命令，在自己的智能体环境安装前先检查源码包。
- 从原始文档核对所需工具、凭据和支持的宿主。
- 说明输入材料与期望交付物，并对照来源检查输出。
npx skills add https://github.com/mohitmishra786/low-level-dev-skills --skill cuda-profiling开始前需了解
- 安装不包含模型订阅、API 额度或访问文件的权限。
- 仓库星数和安装量反映活跃情况，不等于独立质量或安全测试。
仓库原始内容
按源仓库内容呈现，保留标题、案例、代码、表格、链接以及原文引用的演示图片。
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.
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
ncumetrics - 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
Workflow
1. Choose profiling tool
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
bash
# 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-repWhat to look for in the timeline:
- Gaps between kernel launches (CPU bottleneck or sync points)
cudaDeviceSynchronizestalls- Overlap between H2D copies and kernel execution across streams
- CUDA API call overhead
bash
# Capture with CUDA graph info
nsys profile --capture-range=cudaProfilerApi ./my_cuda_app3. NVTX range annotations
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.
4. Nsight Compute — kernel analysis
bash
# 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-repKey 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)
5. NCU CLI metrics
bash
# 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_app6. Memory-bound vs compute-bound diagnosis
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 overheadRoofline 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
7. Occupancy analysis
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 factor | Typical fix |
|---|---|
| Registers | -maxrregcount, simplify kernel |
| Shared memory | Reduce tile size, split phases |
| Block size | Try 128 or 256 instead of 512+ |
8. Profiling workflow checklist
bash
# 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.csvCommon Problems
| Symptom | Cause | Fix |
|---|---|---|
ERR_NVGPUCTRPERM | Insufficient profiling permissions | Run with sudo or set NVreg_RestrictProfilingToAdminUsers=0 |
| All metrics show zero | Profiling disabled or wrong GPU | Check CUDA_VISIBLE_DEVICES; use --target-processes all |
| NCU report empty | Kernel too short or not launched | Increase workload; verify cudaGetLastError() |
| Huge profiling overhead | Full metric sets on many kernels | Use --kernel-name filter; --launch-skip |
| Timeline shows no overlap | Single default stream | Create multiple streams; use async copies |
| Occupancy looks fine but kernel slow | Memory latency not hidden | Check memory coalescing; increase active warps |
Related Skills
skills/gpu/cuda— kernel writing, occupancy tuning, nvcc flagsskills/gpu/gpu-memory-model— coalescing, bank conflicts, SIMT modelskills/gpu/cuda-debugging— correctness before performance tuningskills/profilers/intel-vtune-amd-uprof— CPU-side roofline and hotspot analysisskills/profilers/flamegraphs— CPU flamegraphs complementary to nsys timelineskills/profilers/hardware-counters— general perf stat concepts
