Nsight Compute analysis

Per kernel, with counters

ncu profiles a kernel launch and replays it to collect hardware counters. nsys (gpu-12) records a timeline. Use ncu when the question is why one kernel is slow:

ncu --set full -o <report> ./bench
ncu --section SpeedOfLight --kernel-name regex:mm ./bench

The basic set is the default and is cheap. --set full replays the kernel many times and is not. Section names, metric names, and replay modes change between Nsight Compute releases; the version on the source is where these names came from.

Read the sections in order

SpeedOfLight gives compute and memory throughput as a percentage of peak, which tells you which roof you are under. MemoryWorkloadAnalysis gives sectors per request and hit rates. Occupancy gives achieved warps and the limiter. SchedulerStats and WarpStateStats give issue-slot utilisation and the stall mix. A row of your report looks like this, with your values:

sm__throughput.avg.pct_of_peak_sustained_elapsed        <value>
gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed  <value>
smsp__warps_active.avg.pct_of_peak_sustained_active     <value>

Stalls name a unit

A long-scoreboard stall is a wait on memory. An MIO throttle is a queue at shared memory or the load-store pipe. A barrier stall is __syncthreads(). A wait stall is a fixed-latency dependency. The largest share names the unit to fix, not the fix itself, and if the issue slots are mostly busy then the stall mix is not your problem.

Replay is not free

Multi-pass collection saves and restores memory between passes and serialises launches, so a profiled kernel is slower than the same kernel in a timeline. Never quote an ncu duration as the runtime; quote nsys or CUDA events, and use ncu for attribution.

Things that break

Without -lineinfo the source page has no lines to attribute stalls to. A kernel filter that matches nothing prints nothing and reads as a failed run. On some drivers the counters need root or NVreg_RestrictProfilingToAdminUsers=0, and the error is ERR_NVGPUCTRPERM, which is a permission problem and not a kernel bug.