Profiling Mastery — Nsight + torch.profiler¶

Rule zero: if you don’t have a profile, you don’t have an opinion. This document is the workflow you will run on every kernel you write for the next 12 months.

The three tools — what each is for¶

Tool	Scope	When to use
Nsight Systems (`nsys`)	Whole application timeline: kernels, memcpy, CUDA API, CPU, streams	Find what’s slow across the run — kernel launches, sync stalls, host-device overhead, missed overlap
Nsight Compute (`ncu`)	Single kernel deep dive: SM utilization, memory throughput, warp stalls, instruction mix	Understand why one kernel is slow — optimization is here
`torch.profiler`	PyTorch op-level: which `nn.Linear` / attention / rmsnorm calls, plus CUDA kernels launched by them	Bridge the gap: which PyTorch op → which kernel

These are layered: torch.profiler finds the op, nsys finds the kernel + surrounding timeline, ncu finds the metric.

Latest versions (July 2026)¶

Nsight Compute 2026.2 — https://developer.nvidia.com/tools-overview/nsight-compute/get-started — introduces Nsight Copilot (preview), an LLM-assisted analysis overlay. Worth trying; do not depend on it.
Nsight Systems ships in step with the CUDA Toolkit. Use whatever ships with your CUDA 13.x install.

The universal workflow¶

(1) Have a testable script                              
     │  python bench_my_kernel.py                       
     │                                                  
(2) Whole-run view                                      
     │  nsys profile -o run.qdrep python bench_my_kernel.py
     │  → open in Nsight Systems GUI                     
     │  → find the slow region                          
     │                                                  
(3) Zoom into the specific kernel                       
     │  ncu --set full --launch-skip 100 --launch-count 1 -o kern.ncu-rep python bench_my_kernel.py
     │  → open in Nsight Compute                        
     │  → read speed-of-light + memory workload analysis
     │                                                  
(4) PyTorch context                                     
     │  with torch.profiler.profile(...) as p:          
     │      model(x)                                    
     │  print(p.key_averages().table(sort_by='cuda_time_total'))

Nsight Systems — the timeline view¶

Command:

nsys profile \
  --trace=cuda,nvtx,osrt,cudnn,cublas \
  --cuda-memory-usage=true \
  --output=run \
  python bench_my_kernel.py

Open run.nsys-rep (or older .qdrep) in the Nsight Systems GUI.

What to look for:

Idle gaps on the GPU row → CPU-bound or launch-latency-bound → CUDA graphs will help.
Serial kernel dependency with no overlap → put on different streams / use CUDA graphs.
cudaMemcpy on the critical path → move to pinned memory + async / co-locate.
cuBLAS/cuDNN kernels dominating → good, that’s where you want your time.
Memset/small setup kernels between big kernels → indicates missing fusion.

NVTX ranges are your best friend — wrap logical sections in torch.cuda.nvtx.range_push('rmsnorm') / range_pop() and they appear on the timeline. Do this for every phase of every model.

Nsight Compute — the kernel scalpel¶

Basic command:

ncu --set full \
    --launch-skip 100 --launch-count 1 \
    -o kern \
    python bench_my_kernel.py

--launch-skip 100 --launch-count 1 skips warm-up and profiles one steady-state launch. --set full collects the full metric set (slow; targeted sets like --set roofline or --section MemoryWorkloadAnalysis are faster during iteration).

The report you actually read:

Speed of Light (SOL) section. Two bars: SM% and Memory%.
- Both < 60% → latency-bound (launch, occupancy, sync). Fix stalls first.
- SM% high, Memory% low → compute-bound. Push tensor cores / precision / occupancy.
- Memory% high, SM% low → memory-bound. Fuse, tile, or increase intensity.
Memory Workload Analysis. Traffic through L1/L2/DRAM, hit rates, sector counts. This tells you whether your tiling actually reduced HBM traffic.
Warp State (“Stall Analysis”). Which reason are warps waiting?
- Long Scoreboard → HBM stalls. Add prefetch / pipeline.
- Short Scoreboard → SMEM latency, likely bank conflicts.
- Wait / Barrier → sync overhead, warp imbalance.
- Not Selected → occupancy healthy; scheduler had other warps.
Roofline (2D scatter, per-section). Automatically placed on the machine’s roofline. Fastest way to know where you sit.
Instruction Statistics. Non-matmul FLOP fraction, LDS/STS count. On modern kernels you want tensor-core instructions to dominate.

Metric cheatsheet (metric ID → what it tells you):

Metric	Meaning
`sm__throughput.avg.pct_of_peak_sustained_elapsed`	Overall SM utilization
`sm__pipe_tensor_op_hmma_cycles_active.avg.pct_of_peak_sustained_elapsed`	Tensor core utilization (HMMA-family)
`dram__throughput.avg.pct_of_peak_sustained_elapsed`	HBM utilization
`smsp__average_warps_active.pct_of_peak_sustained_elapsed`	Occupancy
`l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum`	SMEM bank conflicts
`smsp__inst_executed.avg`	Instructions executed
`smsp__thread_inst_executed_per_inst_executed.ratio`	Divergence health (32 = no divergence)

Fixing `ERR_NVGPUCTRPERM` (you will hit this)¶

On shared/cloud GPUs, ncu often fails with ERR_NVGPUCTRPERM — profiler counters not enabled. Two fixes:

Bare-metal: as root, modprobe nvidia NVreg_RestrictProfilingToAdminUsers=0 and reboot, or add to /etc/modprobe.d/.
Cloud/managed: you often cannot fix; use --metrics with only the counters that aren’t restricted (SOL and workload sections mostly work; warp state may not).

The Spheron 2026 guide covers this in detail: https://www.spheron.network/blog/gpu-profiling-ai-workloads-nsight-compute-pytorch-profiler-guide

`torch.profiler` — the PyTorch bridge¶

import torch
from torch.profiler import profile, ProfilerActivity, schedule

def trace_handler(p):
    p.export_chrome_trace(f'trace_{p.step_num}.json')
    print(p.key_averages().table(sort_by='cuda_time_total', row_limit=20))

with profile(
    activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA],
    schedule=schedule(wait=1, warmup=2, active=3, repeat=1),
    on_trace_ready=trace_handler,
    record_shapes=True,
    with_stack=False,          # True for stack traces (slow)
) as p:
    for _ in range(10):
        model(x)
        p.step()

Two things it gives you:

Op-level table: which PyTorch op consumed CUDA time. Great for finding “wait, why is contiguous() taking 8% of my forward pass?”
Chrome trace (drop into chrome://tracing or the Perfetto UI): timeline of ops + kernels.

Pair with NVTX — torch.cuda.nvtx.range() calls appear both in nsys and in the torch trace.

The benchmark hygiene rules (non-negotiable)¶

Warm up before measuring. JIT compilation, cuDNN autotune, kernel selection all happen on first call. Skip ≥ 3 iterations.
Synchronize. torch.cuda.synchronize() before starting the timer and before stopping. Or use CUDA events.
Median, not mean. Report median + p95/p99 for latency-sensitive numbers.
Fix clocks if you care about apples-to-apples: sudo nvidia-smi -lgc 1980 (or whatever base clock).
Isolate the kernel. For ncu specifically, don’t run other CUDA processes on the same GPU.
Verify correctness FIRST. A fast wrong kernel is not fast. torch.allclose(y, ref, atol=..., rtol=...) with the right tolerance for your dtype.
Report the setup. GPU model, driver, CUDA toolkit, dtype, batch size, seqlen, hidden dim. Missing this = incomparable.
Vary shapes. Test at least 3 shape regimes: tiny (batch=1), medium, large. Kernels that shine at large shapes often collapse at small.
Use triton.testing.do_bench for Triton kernels — it handles warmup + sync + statistics correctly.
Log everything. Every benchmark you run should append to a CSV with kernel, shape, dtype, GPU, achieved TFLOPs/s, achieved GB/s, and Nsight report path.

Diagnosis flowchart¶

Kernel is slower than expected.
│
├── Ncu SOL: SM% < 30% AND Mem% < 30%
│     → Latency-bound. Check:
│         - `smsp__average_warps_active` (occupancy)
│         - Launch overhead (nsys)
│         - Sync barriers
│
├── SM% > 70%, Mem% low
│     → Compute-bound. Check:
│         - Tensor-core utilization
│         - Non-matmul FLOP fraction (should be low)
│         - Instruction mix
│
├── Mem% > 70%, SM% low
│     → Memory-bound. Check:
│         - Arithmetic intensity vs ridge
│         - Fuse with neighbors
│         - Tile larger
│         - Check L2 hit rate; maybe your working set doesn't fit
│
└── SM% ~50%, Mem% ~50%
      → Balanced. You're near the roofline. Squeeze last 20% or move on.

Recommended lectures / reads¶

NVIDIA “Intro to Nsight Compute” (2023, still current fundamentals): https://www.youtube.com/watch?v=Iuy_RAvguBM
Kernel Profiling Guide: https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html
Spheron guide (2026, remote/cloud-friendly): https://www.spheron.network/blog/gpu-profiling-ai-workloads-nsight-compute-pytorch-profiler-guide
“Performing Kernel Surgery: Profiling a Matmul Kernel” (practitioner writeup): https://themlsurgeon.substack.com/p/performing-kernel-surgery-profiling
GPU MODE lecture on profiling (search the playlist)
PyTorch profiler docs: https://pytorch.org/docs/stable/profiler.html

What “done” looks like¶

You can profile any kernel you write and produce a one-page summary: SOL numbers, top-3 stall reasons, roofline position, top fix.
You can navigate the nsys timeline and find kernel launch gaps.
You have NVTX ranges around every named phase in your inference stack.
You have a CSV log of every benchmark you’ve run since Phase 2 started. This becomes your portfolio evidence.