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

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

# 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

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

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

3. NVTX range annotations

#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

# 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

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)

5. NCU CLI metrics

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

7. Occupancy analysis

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

# 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

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

cuda-profiling