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

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

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

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

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

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

NPX Install

Tags

SKILL.md Content

CUDA Profiling

Purpose

When to Use

Workflow

1. Choose profiling tool

2. Nsight Systems — timeline profiling

3. NVTX range annotations

4. Nsight Compute — kernel analysis

5. NCU CLI metrics

6. Memory-bound vs compute-bound diagnosis

7. Occupancy analysis

8. Profiling workflow checklist

Common Problems

Related Skills