Nsight Systems & Nsight Compute

NVIDIA's profilers, from an infrastructure point of view: what each Nsight tool answers, how to run nsys and ncu on servers, containers, and Slurm clusters, how to read a timeline and a kernel report, how they compare with nvidia-smi, DCGM, and nccl-tests, and what breaks (permissions, counter conflicts, overhead).

1. The big picture: the Nsight family

"Nsight" is a family of NVIDIA developer tools. For AI infrastructure, two matter:

Tool	CLI	Question it answers	Scope
Nsight Systems	`nsys`	"Where does the time go across my whole application?"	System-wide timeline: CPU threads, CUDA API calls, GPU kernels, memory copies, NCCL, OS calls, MPI, NVTX ranges, across one or many processes and nodes
Nsight Compute	`ncu`	"Why is this kernel slow?"	Deep dive into individual GPU kernels: throughput, memory traffic, occupancy, stalls, roofline

Others you may see named: Nsight Graphics (graphics frame debugging), Nsight Deep Learning Designer, and IDE integrations (Nsight Visual Studio Edition, Nsight Visual Studio Code Edition, Nsight Eclipse Edition) for debugging CUDA code. They're developer tools, not cluster operations tools.

Nsight Systems is a traffic camera over the whole city: you see where cars queue, which roads are empty, and when the bridge is jammed. Nsight Compute is a mechanic's diagnostic on one car: it tells you exactly why that car is slow, but only for the car you put on the lift.

Work top-down: start with Nsight Systems to find where the time goes (idle GPUs, data loading, communication, one hot kernel). Only then use Nsight Compute on the specific kernels that dominate. Running ncu on a whole training job first is slow and usually answers the wrong question.

Systems = Schedule (timeline). Compute = Close-up of one kernel.

2. Where Nsight fits among infra tools

Tool	Granularity	Typical user	Use it for
`nvidia-smi`	Per GPU, a snapshot every second or so	Admin	Is the GPU present, busy, hot, full of memory? (guide)
DCGM / dcgm-exporter	Per GPU, continuous, fleet-wide	Admin / SRE	Health, diagnostics, monitoring metrics such as SM activity and tensor-core activity for every job (guide)
nccl-tests	Collective bandwidth	Admin	Is the fabric delivering expected bandwidth? (guide)
Nsight Systems	Per event (microseconds), one job, short window	Developer, performance engineer, admin debugging a slow job	Why is this job slow? GPU gaps, data loading, communication vs compute overlap, stragglers
Nsight Compute	Per kernel, hardware counters	Kernel / framework developer	Why is this kernel slow? Memory-bound vs compute-bound, occupancy

nvidia-smi "GPU-Util 100%" only means some kernel was running during the sample period. It doesn't say the GPU was used efficiently. DCGM's SM activity / occupancy metrics, and then Nsight, tell you how much of the GPU was actually doing work.

3. Nsight Systems: what it records

Trace source	What you see
`cuda`	CUDA API calls on the CPU (launches, `cudaMemcpy`, syncs) and the matching GPU kernels and copies on each GPU stream
`nvtx`	Named ranges added by the application or framework (for example "forward", "backward", "optimizer step"). NCCL and many frameworks emit them too.
`osrt`	OS runtime calls: file reads, `pthread` waits, sleeps (data-loading stalls show up here)
`cublas`, `cudnn`	Library calls
`mpi`, `ucx`	MPI and UCX communication calls
CPU sampling	Which CPU functions were running (backtraces)
GPU metrics sampling	Periodic hardware metrics per GPU: SM activity, tensor activity, memory bandwidth, PCIe and NVLink throughput
NIC metrics	InfiniBand / Ethernet adapter throughput on supported ConnectX NICs (newer versions)

Overhead is usually low (a few percent), so nsys is suitable for profiling realistic training steps. Keep the window short: reports grow quickly.

4. nsys profile: the flags that matter

nsys profile \
  -o train_profile \                        # output: train_profile.nsys-rep
  --trace=cuda,nvtx,osrt,cudnn,cublas \     # what to trace
  --sample=cpu \                            # CPU sampling (or none to reduce overhead)
  --stats=true \                            # print summary tables at the end
  --delay=60 --duration=30 \                # skip warm-up, record 30 s
  --force-overwrite=true \
  python train.py

Flag	Meaning
`-o` / `--output`	Report name. Supports substitutions such as `%h` (hostname), `%p` (process ID), and `%q{VAR}` (an environment variable, for example the Slurm rank).
`-t` / `--trace`	Comma-separated trace sources: `cuda,nvtx,osrt,cublas,cudnn,mpi,ucx,...`
`-s` / `--sample`	CPU sampling: `cpu` or `none`
`--stats=true`	Print summary reports when profiling ends
`-y` / `--delay`, `-d` / `--duration`	Start after N seconds; stop after N seconds
`-c` / `--capture-range`	Only record a region: `cudaProfilerApi` (between `cudaProfilerStart/Stop`) or `nvtx` (inside a named NVTX range, with `--nvtx-capture`)
`--gpu-metrics-devices`	Enable GPU metrics sampling on selected GPUs (for example `all`)
`--cuda-memory-usage=true`	Track GPU memory allocations over time
`--nic-metrics=true`	Sample NIC counters on supported adapters
`-f` / `--force-overwrite`	Overwrite an existing report

Flag names have changed across Nsight Systems releases (for example the GPU metrics flag was singular in older versions, and older reports used the .qdrep extension instead of .nsys-rep). Run nsys profile --help for your installed version.

Interactive sessions

nsys launch --trace=cuda,nvtx python serve.py   # start the app under nsys, not recording yet
nsys start  -o serve_peak                        # begin recording when load arrives
nsys stop                                        # finish and write the report
nsys sessions list

Useful for long-running services (inference servers) where you only want to capture a busy moment.

5. Reports: nsys stats, export, and the GUI

nsys stats train_profile.nsys-rep                          # all default summary tables
nsys stats --report cuda_gpu_kern_sum train_profile.nsys-rep   # top GPU kernels by total time
nsys stats --report cuda_api_sum      train_profile.nsys-rep   # CUDA API calls (syncs, launches)
nsys stats --report cuda_gpu_mem_time_sum train_profile.nsys-rep  # memcpy / memset time
nsys stats --report nvtx_sum          train_profile.nsys-rep   # time per NVTX range
nsys stats --report osrt_sum          train_profile.nsys-rep   # OS calls (I/O, waits)
nsys export --type sqlite train_profile.nsys-rep           # for custom SQL / scripts
nsys analyze train_profile.nsys-rep                        # rule-based hints (where available)

The usual workflow on clusters: record on the compute node with the CLI (no GUI needed), copy the .nsys-rep file to a workstation, and open it in the Nsight Systems GUI to view the timeline. The GUI can also open several reports together (for example one per rank) to compare them.

Report names in nsys stats have been renamed between versions. nsys stats --help-reports lists the ones your version supports.

6. Reading a timeline

This timeline shows a classic input-bound job: the GPU rows have large gaps, and each gap lines up with the CPU data loader being busy. Faster GPUs or networks won't help; faster data loading will.

Pattern on the timeline	Likely meaning	Typical fix
Gaps in GPU rows while CPU threads read files or decode data (`osrt` shows `read`, waits)	Input pipeline / storage bound	More loader workers, faster storage or local cache (storage guide), prefetching, GPU data loading (DALI)
Many tiny kernels with gaps between them; CPU busy launching	CPU launch overhead	Larger batches, kernel fusion, CUDA Graphs
Long `cudaMemcpy` HtoD/DtoH rows, marked pageable	Copies from non-pinned memory	Pinned (page-locked) host memory, asynchronous copies overlapped with compute
Frequent `cudaDeviceSynchronize` / `cudaStreamSynchronize`	Unneeded syncs stall the pipeline (for example `.item()` in Python every step)	Remove syncs from the hot loop
NCCL kernels long and not overlapped with compute	Communication-bound, or overlap disabled	Check fabric with nccl-tests; enable gradient bucketing / overlap
NCCL kernels long on all ranks except one, which arrives late	A straggler: other ranks wait inside the collective for the slow rank	Compare ranks; check the slow node (thermals, clocks, data, NIC)
GPU rows dense, one kernel dominates `cuda_gpu_kern_sum`	Genuinely compute-bound on that kernel	Now use Nsight Compute on that kernel

A long NCCL kernel doesn't automatically mean a slow network. In a collective, every rank waits for the slowest one, so a fast rank shows a long AllReduce while it waits for a straggler. Compare the start times of the collective across ranks before blaming the fabric.

7. NVTX: labeling your code

NVTX (NVIDIA Tools Extension) is a small API for marking named ranges and events in code. Nsight Systems shows them as labeled bars, which turns an anonymous wall of kernels into "data load → forward → backward → AllReduce → optimizer".

# PyTorch
import torch
torch.cuda.nvtx.range_push("forward")
out = model(batch)
torch.cuda.nvtx.range_pop()

# Python nvtx package
import nvtx
with nvtx.annotate("data_load", color="orange"):
    batch = next(loader)

NCCL emits its own NVTX ranges for collectives, so with --trace=nvtx you can see which collective (and how many bytes) each NCCL kernel belongs to. NVTX ranges also power --capture-range=nvtx to record only a chosen region.

8. Multi-GPU and multi-node profiling

Put nsys between the launcher and the application, so every rank is profiled and writes its own report:

# Slurm: one report per rank, named by hostname and rank
srun nsys profile -o prof_%h_%q{SLURM_PROCID} --trace=cuda,nvtx,osrt,mpi \
     --delay=120 --duration=20 python train.py

# MPI
mpirun -np 16 nsys profile -o prof_%q{OMPI_COMM_WORLD_RANK} --trace=cuda,nvtx,mpi ./app

# torchrun: profile the launcher and its child processes
nsys profile -o prof_%h --trace=cuda,nvtx,osrt torchrun --nproc_per_node=8 train.py

Don't wrap srun itself (nsys profile srun ...): that profiles the launcher on one node, not the ranks.
Write reports to a shared or node-local scratch path, and keep durations short: 64 ranks × 30 seconds of tracing produces a lot of data.
Profiling a subset of ranks (for example one per node) is often enough to find stragglers and keeps the data manageable.
Open several rank reports together in the GUI to line up collectives across ranks.

For slow multi-node training, Nsight Systems answers "is it compute, data, or communication, and is one rank holding everyone back?". If communication is slow on every rank, validate the fabric next with nccl-tests and the InfiniBand diagnostics.

9. Nsight Compute: how it works

Kernel by kernel: ncu intercepts each selected kernel launch, runs it while collecting hardware performance counters, and records a detailed report.
Replay: only a limited number of counters can be collected at once, so ncu replays the same kernel several times, saving and restoring GPU memory between passes, to collect everything requested. This is why it's slow.
Serialization: profiled kernels run one at a time, so concurrency between streams isn't representative.
Clock control: by default ncu locks GPU clocks to base clock, so results are repeatable between runs. That's also why profiled kernels run slower than in production. --clock-control none leaves clocks alone.
Sets and sections: metrics are grouped into sections (for example SpeedOfLight, MemoryWorkloadAnalysis, Occupancy), and sections into sets (a small default set, or full).

Nsight Compute adds large overhead (each profiled kernel can take many times longer), and changes clocks and concurrency. It's a lab tool for a few kernels, not something to attach to a whole production training run. For whole-job timing, use Nsight Systems.

10. ncu: the flags that matter

ncu -o gemm_report \                 # output: gemm_report.ncu-rep
    --set full \                      # collect all sections (slow); default set is smaller
    -k regex:gemm \                   # only kernels whose name matches
    -s 50 -c 5 \                      # skip the first 50 matching launches, profile 5
    python bench.py

ncu --list-sets                       # available sets
ncu --list-sections                   # available sections
ncu --section SpeedOfLight --section Occupancy ./app
ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed ./app
ncu --target-processes all python train.py   # include child processes
ncu --import gemm_report.ncu-rep --page details   # print a saved report in the terminal

Flag	Meaning
`-o`	Save to a `.ncu-rep` report (open in the Nsight Compute GUI)
`--set`	Which group of sections to collect (`full` = everything)
`--section` / `--metrics`	Collect specific sections or individual metrics
`-k` / `--kernel-name`	Filter kernels by name (exact, or `regex:`)
`-s` / `--launch-skip`, `-c` / `--launch-count`	Skip warm-up launches; limit how many are profiled
`--replay-mode`	`kernel` (default), `application` (rerun the whole app per pass), `range`
`--clock-control`	`base` (default, locked) or `none`
`--target-processes all`	Follow child processes (Python launchers, MPI)

11. Reading a kernel report

Section	What it tells you
GPU Speed Of Light (SOL)	Achieved compute (SM) throughput and memory throughput as a percentage of the GPU's peak. The first thing to read.
Roofline	Plots the kernel's arithmetic intensity (FLOPs per byte) against the hardware limits: left of the ridge = memory-bound, right = compute-bound
Memory Workload Analysis	Traffic through L1, L2, and HBM; cache hit rates
Occupancy	Active warps per SM vs the maximum, and what limits it (registers, shared memory, block size)
Launch Statistics	Grid and block size, registers per thread. Tiny grids can't fill a large GPU.
Warp State / Scheduler Statistics	Why warps stall (waiting on memory, dependencies, synchronization)

SOL result	Interpretation
Memory % high, compute % low	Memory-bound: improve data reuse, fusion, lower precision
Compute % high, memory % lower	Compute-bound: use tensor cores / lower precision, better algorithms
Both low	Latency-bound or under-occupied: too few threads, small grids, stalls; check launch stats and occupancy

These are developer-level interpretations. For an infrastructure role, the important point is recognizing when a problem is "inside the kernel" (hand it to developers with an ncu report) versus "outside the kernel" (data, communication, configuration, hardware), which Nsight Systems, DCGM, and the fabric tools address.

12. Permissions, containers, and environment

GPU performance counter permissions

By default, the NVIDIA driver restricts access to GPU performance counters to administrators. A normal user running ncu (or nsys GPU metrics) gets an error like ERR_NVGPUCTRPERM: "The user does not have permission to access NVIDIA GPU Performance Counters".

# Option 1: run the profiler with admin rights (sudo / root)

# Option 2: allow all users (security trade-off: counters can reveal information about other users' work)
echo 'options nvidia NVreg_RestrictProfilingToAdminUsers=0' | sudo tee /etc/modprobe.d/nvidia-profiling.conf
# then reload the driver or reboot (and rebuild the initramfs if the driver loads early)
cat /proc/driver/nvidia/params | grep RmProfilingAdminOnly   # 0 = unrestricted

Counter access is a driver-level setting on the host. Installing a newer Nsight version doesn't fix ERR_NVGPUCTRPERM; admin rights or the NVreg_RestrictProfilingToAdminUsers module parameter does. Plain CUDA tracing with nsys generally doesn't need counter access.

Containers

NGC framework containers (PyTorch, TensorFlow) usually include nsys and often ncu; otherwise mount or install them.
ncu in a container typically needs extra privileges (for example --cap-add=SYS_ADMIN) in addition to the host's counter-permission setting.
CPU sampling in nsys depends on the kernel's perf_event_paranoid setting (lower = more permissive), which containers inherit from the host.

nsys status --environment            # checks whether CPU sampling, tracing, etc. will work here
cat /proc/sys/kernel/perf_event_paranoid
docker run --gpus all --cap-add=SYS_ADMIN -v $PWD:/work nvcr.io/nvidia/pytorch:<tag> \
    ncu --set full -k regex:gemm -c 3 python /work/bench.py

Slurm clusters

Administrators often install Nsight on shared storage or as an environment module, and decide cluster-wide whether profiling counters are restricted. Some sites expose profiling through a job option or a dedicated partition rather than opening counters to everyone.

13. Counter conflicts and overhead

Hardware performance counters are a shared, limited resource on each GPU. Only one tool can collect profiling metrics at a time:

DCGM profiling metrics (the "DCP" fields such as SM activity and tensor activity, used by dcgm-exporter) hold the counters continuously.
Nsight Compute and Nsight Systems GPU metrics sampling need them too.

dcgmi profile --pause      # release the counters for this node's GPUs
ncu ... / nsys profile --gpu-metrics-devices=all ...
dcgmi profile --resume     # give them back to DCGM monitoring

Tool	Typical overhead	Safe on production jobs?
`nsys` (CUDA + NVTX trace)	Low (a few percent); higher with heavy CPU sampling or many tiny kernels	Yes, for short windows
`ncu`	High: each profiled kernel replays many times; clocks locked	No; profile specific kernels in a test run
DCGM profiling metrics	Very low	Yes, continuously

If ncu, nsys GPU metrics, or dcgmproftester fails with a profiling or "counters in use" error on a monitored cluster, think DCGM holds the counters: pause DCGM profiling, profile, then resume. See the DCGM deep dive and gpu-burn & dcgmproftester.

14. Nsight vs other tools

Question	Best tool	Not this
Is the GPU present, hot, or out of memory?	`nvidia-smi`	Nsight
How busy were all GPUs in the cluster last week?	DCGM / dcgm-exporter + Prometheus	Nsight (one job, short window)
Is the hardware healthy?	`dcgmi diag`, NVSM	Nsight
Does the fabric reach expected AllReduce bandwidth?	nccl-tests	Nsight (it shows a job's communication, but doesn't benchmark the fabric)
Why does this training step take 800 ms?	Nsight Systems	Nsight Compute (too narrow), nvidia-smi (too coarse)
Why is this one GEMM kernel at 30% of peak?	Nsight Compute	Nsight Systems (shows how long, not why)
Python-level operator breakdown in PyTorch	PyTorch Profiler (can also emit NVTX for Nsight)	—

15. Command builder

Goal
Launcher
Application

16. "Which tool?" chooser

Pick a question above.

17. Sample outputs

What the commands in this guide print. Expand a command to see a typical result and what to check in it.

Illustrative output, not captured from a real system. Hostnames, serial numbers, GUIDs, MAC and IP addresses are invented, and columns and wording vary between versions. Learn the shape of each output and the fields called out under it.

nsys profile -o train_profile --trace=cuda,nvtx,osrt,cudnn,cublas --sample=cpu --stats=true --delay=60 --duration=30 --force-overwrite=true python train.py

Record a timeline with summary tables

WARNING: CPU IP/backtrace sampling not supported, disabling.
Try the 'nsys status --environment' command to learn more.

Collection is delayed for 60 seconds.
epoch 0 step 100 loss 2.913 (app output continues)
Collecting data...
epoch 0 step 400 loss 2.402
Generating '/tmp/nsys-report-3f1a.qdstrm'
[1/8] [========================100%] train_profile.nsys-rep
[2/8] [========================100%] train_profile.sqlite
[3/8] Executing 'nvtx_sum' stats report
...
[5/8] Executing 'cuda_api_sum' stats report
...
[6/8] Executing 'cuda_gpu_kern_sum' stats report
...
[8/8] Executing 'cuda_gpu_mem_size_sum' stats report
...
Generated:
    /home/alice/train_profile.nsys-rep
    /home/alice/train_profile.sqlite

What to look at: Confirm the Generated: lines list a .nsys-rep and that 'Collecting data...' appeared after the delay. The CPU sampling warning means perf_event_paranoid is too strict or the container lacks privileges; tracing still works but you lose CPU call stacks.

nsys stopFinish an interactive nsys start capture

Generating '/tmp/nsys-report-8c2d.qdstrm'
[1/1] [========================100%] serve_peak.nsys-rep
Generated:
    /home/alice/serve_peak.nsys-rep

What to look at: After nsys launch and nsys start -o serve_peak, stop writes the report for just the recorded window. 'No active sessions' means the start was run as another user or in another shell session than the launch.

nsys stats --report cuda_gpu_kern_sum train_profile.nsys-repTop GPU kernels by total time

Processing [train_profile.sqlite] with [/usr/local/cuda-12.4/nsight-systems-2023.4.4/host-linux-x64/reports/cuda_gpu_kern_sum.py]...

 ** CUDA GPU Kernel Summary (cuda_gpu_kern_sum):

 Time (%)  Total Time (ns)  Instances   Avg (ns)     Med (ns)    Min (ns)   Max (ns)   StdDev (ns)                     Name
 --------  ---------------  ---------  -----------  -----------  ---------  ---------  -----------  ------------------------------------------------
     38.2   11,204,118,330      6,120  1,830,738.3  1,829,120.0  1,801,344  1,902,912     12,004.7  sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x256x64_…
     17.9    5,250,334,912      3,060  1,715,795.7  1,714,560.0  1,690,208  1,770,016      9,881.2  flash_fwd_kernel<Flash_fwd_kernel_traits<128, 128, 64, 4…
     11.4    3,343,982,104      1,530  2,185,609.2  2,184,192.0  2,150,912  2,240,128     14,550.3  ncclDevKernel_AllReduce_Sum_bf16_RING_LL128(ncclDevKernel…
      6.1    1,789,227,615     24,480     73,089.4     72,896.0     70,560     79,872      1,022.7  void at::native::vectorized_elementwise_kernel<4, at::nat…
...

What to look at: Kernels are sorted by Time (%); the top few are where GPU time goes. A large share in NCCL kernels points at communication or a straggler rank, and thousands of tiny kernels with small Avg (ns) suggest launch overhead (try CUDA Graphs).

nsys stats --report cuda_api_sum train_profile.nsys-repCUDA API calls: syncs and launches

Processing [train_profile.sqlite] with [/usr/local/cuda-12.4/nsight-systems-2023.4.4/host-linux-x64/reports/cuda_api_sum.py]...

 ** CUDA API Summary (cuda_api_sum):

 Time (%)  Total Time (ns)  Num Calls    Avg (ns)      Med (ns)     Min (ns)   Max (ns)     StdDev (ns)           Name
 --------  ---------------  ---------  ------------  ------------  --------  ------------  -----------  ---------------------
     61.4   18,402,113,876      1,830  10,055,800.0   9,812,441.0     2,211   211,004,118  7,120,554.3  cudaStreamSynchronize
     21.7    6,503,881,201    412,330      15,773.5       6,104.0     3,018    12,883,410     71,220.9  cudaLaunchKernel
      9.8    2,937,004,512      3,672     799,837.8     512,330.0     8,102    18,220,441    990,334.1  cudaMemcpyAsync
      3.1      929,118,004        612   1,518,166.7   1,503,220.0   402,118     4,120,330    311,022.6  cudaHostAlloc
...

What to look at: This is CPU-side time spent inside CUDA calls. A dominant cudaStreamSynchronize or cudaDeviceSynchronize means the host is waiting on the GPU (often fine), while high cudaLaunchKernel time with many calls means the CPU cannot launch kernels fast enough.

nsys stats --report cuda_gpu_mem_time_sum train_profile.nsys-repMemcpy and memset time on the GPU

Processing [train_profile.sqlite] with [/usr/local/cuda-12.4/nsight-systems-2023.4.4/host-linux-x64/reports/cuda_gpu_mem_time_sum.py]...

 ** CUDA GPU MemOps Summary (by Time) (cuda_gpu_mem_time_sum):

 Time (%)  Total Time (ns)  Count   Avg (ns)    Med (ns)   Min (ns)   Max (ns)   StdDev (ns)            Operation
 --------  ---------------  -----  ----------  ----------  --------  ----------  -----------  ----------------------------
     82.6    1,402,330,118  1,836   763,796.4   761,120.0   740,192   1,120,448     21,004.1  [CUDA memcpy Host-to-Device]
     14.1      239,551,004  1,224   195,711.6   194,880.0     2,112     402,336     98,330.7  [CUDA memcpy Device-to-Host]
      3.3       56,021,887  2,448    22,884.8    22,656.0     1,184      48,288      9,771.2  [CUDA memset]

What to look at: Host-to-Device time that rivals kernel time means the input pipeline or pageable (non-pinned) memory is the bottleneck. Check that memcpy Avg (ns) is consistent; a long tail in Max (ns) often lines up with GPU idle gaps on the timeline.

nsys stats --report nvtx_sum train_profile.nsys-repTime per NVTX range

Processing [train_profile.sqlite] with [/usr/local/cuda-12.4/nsight-systems-2023.4.4/host-linux-x64/reports/nvtx_sum.py]...

 ** NVTX Range Summary (nvtx_sum):

 Time (%)  Total Time (ns)  Instances    Avg (ns)       Med (ns)      Min (ns)     Max (ns)     StdDev (ns)    Style      Range
 --------  ---------------  ---------  -------------  -------------  -----------  -----------  ------------  -------  ---------
     52.3   15,690,221,004        306   51,275,232.0   51,102,448.0  49,880,112   58,330,991   1,102,334.7  PushPop  backward
     27.9    8,370,118,442        306   27,353,328.2   27,220,110.0  26,901,442   31,004,887     611,220.4  PushPop  forward
     19.8    5,940,330,118        306   19,412,843.5    2,104,330.0   1,880,221  402,118,337  61,330,118.9  PushPop  data_load

What to look at: NVTX ranges map GPU work back to code phases. Here data_load has a median of about 2 ms but a max of 402 ms and a huge StdDev, which is the signature of an input-bound job stalling on storage or the dataloader.

nsys stats --report osrt_sum train_profile.nsys-repOS runtime calls: I/O and waits

Processing [train_profile.sqlite] with [/usr/local/cuda-12.4/nsight-systems-2023.4.4/host-linux-x64/reports/osrt_sum.py]...

 ** OS Runtime Summary (osrt_sum):

 Time (%)  Total Time (ns)  Num Calls     Avg (ns)        Med (ns)      Min (ns)     Max (ns)       StdDev (ns)          Name
 --------  ---------------  ---------  --------------  -------------  ---------  -------------  ---------------  ----------------
     58.1  121,330,442,118      4,412    27,500,100.2   10,001,224.0      1,220  2,003,118,004    112,004,330.1  poll
     30.2   63,071,004,887     18,330     3,440,862.2      1,902,118.0      2,104    402,118,220     14,220,118.6  pthread_cond_wait
      8.4   17,544,118,330    212,004        82,753.9         21,330.0      1,004     88,330,442        401,220.9  read
...

What to look at: poll and pthread_cond_wait are usually idle helper threads and are expected to be large. Long read or open calls on the training threads, especially with big Max (ns), point at slow storage rather than the GPU.

nsys analyze train_profile.nsys-repRule-based hints from a report

Processing [train_profile.sqlite] with [/usr/local/cuda-12.4/nsight-systems-2023.4.4/host-linux-x64/rules/cuda_memcpy_sync.py]...

 ** CUDA Synchronous Memcpy (cuda_memcpy_sync):

There were no problems detected related to synchronous memcpy operations.

Processing [train_profile.sqlite] with [/usr/local/cuda-12.4/nsight-systems-2023.4.4/host-linux-x64/rules/gpu_gaps.py]...

 ** GPU Gaps (gpu_gaps):

The following are ranges where a GPU is idle for more than 500ms. Addressing these
gaps might improve application performance.

 Row#  Duration (ns)    Start (ns)       PID    Device ID  Context ID
 ----  -------------  ---------------  -------  ---------  ----------
    1    402,331,118   72,118,330,442  52114           0           1
    2    388,004,220   91,402,118,004  52114           0           1
...

What to look at: Each rule either reports no problems or lists suspicious ranges. GPU gaps of hundreds of milliseconds that line up with data_load NVTX ranges confirm an input-bound job; synchronous memcpy findings point at pageable host memory.

srun nsys profile -o prof_%h_%q{SLURM_PROCID} --trace=cuda,nvtx,osrt,mpi --delay=120 --duration=20 python train.pyOne report per Slurm rank

Collection is delayed for 120 seconds.
Collection is delayed for 120 seconds.
...
Generating '/tmp/nsys-report-a1b2.qdstrm'
Generating '/tmp/nsys-report-c3d4.qdstrm'
[1/1] [========================100%] prof_dgx01_0.nsys-rep
[1/1] [========================100%] prof_dgx01_1.nsys-rep
...
Generated:
    /lustre/home/alice/prof_dgx01_0.nsys-rep
Generated:
    /lustre/home/alice/prof_dgx02_8.nsys-rep
...

What to look at: nsys goes after srun so every rank gets its own profiler, and %h / %q{SLURM_PROCID} make the names unique. If all ranks write the same file name (for example prof.nsys-rep), reports overwrite each other and you are left with one random rank.

ncu -o gemm_report --set full -k regex:gemm -s 50 -c 5 python bench.pyDeep-profile five matching GEMM kernels

==PROF== Connected to process 52114 (/usr/bin/python3.10)
==PROF== Profiling "sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesiz..." - 0: 0%....50%....100% - 41 passes
==PROF== Profiling "sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesiz..." - 1: 0%....50%....100% - 41 passes
==PROF== Profiling "sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesiz..." - 2: 0%....50%....100% - 41 passes
==PROF== Profiling "sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesiz..." - 3: 0%....50%....100% - 41 passes
==PROF== Profiling "sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesiz..." - 4: 0%....50%....100% - 41 passes
==PROF== Disconnected from process 52114
==PROF== Report: /home/alice/gemm_report.ncu-rep

What to look at: Each profiled kernel is replayed many times (41 passes for the full set), which is why ncu is slow and must be filtered with -k, -s and -c. '==WARNING== No kernels were profiled' means the regex or skip count matched nothing.

ncu --list-setsAvailable section sets

---------- --------------------------------------------------------------------------- ------- -----------------
Identifier Sections                                                                    Enabled Estimated Metrics
---------- --------------------------------------------------------------------------- ------- -----------------
basic      LaunchStats, Occupancy, SpeedOfLight, WorkloadDistribution                  yes     190
detailed   ComputeWorkloadAnalysis, LaunchStats, MemoryWorkloadAnalysis, MemoryWorkloa no      619
           dAnalysis_Chart, Occupancy, SourceCounters, SpeedOfLight, SpeedOfLight_Roof
           lineChart, WorkloadDistribution
full       ComputeWorkloadAnalysis, InstructionStats, LaunchStats, MemoryWorkloadAnaly no      5902
           sis, MemoryWorkloadAnalysis_Chart, MemoryWorkloadAnalysis_Tables, NumaAffin
           ity, Nvlink_Tables, Nvlink_Topology, Occupancy, SchedulerStats, SourceCount
           ers, SpeedOfLight, SpeedOfLight_RooflineChart, WarpStateStats, WorkloadDist
           ribution
nvlink     Nvlink, Nvlink_Tables, Nvlink_Topology                                      no      52
roofline   SpeedOfLight, SpeedOfLight_HierarchicalDoubleRooflineChart, SpeedOfLight_Hi no      5260
           erarchicalHalfRooflineChart, SpeedOfLight_HierarchicalSingleRooflineChart,
           SpeedOfLight_HierarchicalTensorRooflineChart, SpeedOfLight_RooflineChart
source     SourceCounters                                                              no      74

What to look at: basic is the default set (Enabled yes); full collects thousands of metrics and needs many more replay passes. Use the Estimated Metrics column to judge overhead before profiling a long job with --set full.

ncu --section SpeedOfLight --section Occupancy ./appSpeed-of-light and occupancy for a kernel

==PROF== Connected to process 61021 (/home/alice/app)
==PROF== Profiling "vectorAdd" - 0: 0%....50%....100% - 9 passes
==PROF== Disconnected from process 61021
[61021] app@127.0.0.1
  vectorAdd(const float *, const float *, float *, int) (65536, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 9.0
    Section: GPU Speed Of Light Throughput
    ----------------------- ------------- ------------
    Metric Name               Metric Unit Metric Value
    ----------------------- ------------- ------------
    DRAM Frequency          cycle/nsecond         2.62
    SM Frequency            cycle/nsecond         1.75
    Elapsed Cycles                  cycle       60,212
    Memory Throughput                   %        91.47
    DRAM Throughput                     %        91.47
    Duration                      usecond        34.40
    L1/TEX Cache Throughput             %        14.83
    L2 Cache Throughput                 %        38.21
    SM Active Cycles                cycle    52,118.62
    Compute (SM) Throughput             %        16.35
    ----------------------- ------------- ------------

    INF   The kernel is utilizing greater than 80.0% of the available compute or memory performance of the device.
          To further improve performance, work will likely need to be shifted from the most utilized to another
          unit. Start by analyzing DRAM in the Memory Workload Analysis section.

    Section: Occupancy
    ------------------------------- ----------- ------------
    Metric Name                     Metric Unit Metric Value
    ------------------------------- ----------- ------------
    Block Limit Registers                 block            8
    Block Limit Warps                     block            8
    Theoretical Active Warps per SM        warp           64
    Theoretical Occupancy                     %          100
    Achieved Occupancy                        %        82.14
    Achieved Active Warps Per SM           warp        52.57
    ------------------------------- ----------- ------------

What to look at: Compare Memory Throughput with Compute (SM) Throughput: here memory is at 91% and compute at 16%, so the kernel is memory-bound and is already near the hardware limit. Low values for both usually mean too little work or poor occupancy, which the Achieved Occupancy row shows.

ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed ./appCollect one specific metric

==PROF== Connected to process 61188 (/home/alice/app)
==PROF== Profiling "vectorAdd" - 0: 0%....50%....100% - 1 pass
==PROF== Disconnected from process 61188
[61188] app@127.0.0.1
  vectorAdd(const float *, const float *, float *, int) (65536, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 9.0
    Section: Command line profiler metrics
    ---------------------------------------------- ----------- ------------
    Metric Name                                    Metric Unit Metric Value
    ---------------------------------------------- ----------- ------------
    sm__throughput.avg.pct_of_peak_sustained_elapsed           %        16.35
    ---------------------------------------------- ----------- ------------

What to look at: Collecting a single metric needs only one pass, so overhead is low. The value is a percentage of the SM's sustained peak; a low number on a kernel you expected to be compute-bound means it is limited by memory, latency or launch configuration.

ncu --import gemm_report.ncu-rep --page detailsPrint a saved Nsight Compute report

[52114] python3.10@127.0.0.1
  sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x256x64_warpgroupsiz... (132, 1, 1)x(384, 1, 1), Context 1, Stream 7, Device 0, CC 9.0
    Section: GPU Speed Of Light Throughput
    ----------------------- ------------- ------------
    Metric Name               Metric Unit Metric Value
    ----------------------- ------------- ------------
    DRAM Frequency          cycle/nsecond         2.62
    SM Frequency            cycle/nsecond         1.98
    Elapsed Cycles                  cycle    3,612,118
    Memory Throughput                   %        42.18
    DRAM Throughput                     %        42.18
    Duration                      msecond         1.82
    L1/TEX Cache Throughput             %        37.60
    L2 Cache Throughput                 %        40.92
    SM Active Cycles                cycle 3,598,224.11
    Compute (SM) Throughput             %        78.64
    ----------------------- ------------- ------------

    INF   Compute is more heavily utilized than Memory: Look at the Compute Workload Analysis section to see what
          the compute pipelines are spending their time doing.
...
    Section: Launch Statistics
    Block Size                                                 384
    Grid Size                                                  132
    Registers Per Thread                   register/thread     168
...

What to look at: --import reads a .ncu-rep without re-running the app, so it works on a login node without a GPU. For a healthy GEMM expect Compute (SM) Throughput well above Memory Throughput; a GEMM that is memory-bound or at low SM throughput suggests bad shapes, wrong precision or throttled clocks.

ncu ./appPermission error on GPU counters

==PROF== Connected to process 63304 (/home/alice/app)
==ERROR== ERR_NVGPUCTRPERM - The user does not have permission to access NVIDIA GPU Performance Counters on the target device 0. For instructions on enabling permissions and to get more information see https://developer.nvidia.com/ERR_NVGPUCTRPERM
==PROF== Disconnected from process 63304
==WARNING== No kernels were profiled.

What to look at: ERR_NVGPUCTRPERM means a non-root user tried to read performance counters while the driver restricts them to admins. Fix it by running as root, or by setting NVreg_RestrictProfilingToAdminUsers=0 and reloading the driver; in containers also add --cap-add=SYS_ADMIN.

cat /proc/driver/nvidia/params | grep RmProfilingAdminOnlyIs profiling restricted to admins?

RmProfilingAdminOnly: 1

What to look at: 1 (the default) means only root or CAP_SYS_ADMIN can use ncu and nsys GPU metrics, which explains ERR_NVGPUCTRPERM for normal users. 0 means the NVreg_RestrictProfilingToAdminUsers=0 option took effect after the driver reload.

nsys status --environmentCheck what nsys can collect here

Timestamp counter supported: Yes

CPU Profiling Environment Check
Root privilege: disabled
Linux Kernel Paranoid Level = 4
Linux Distribution = Ubuntu
Linux Kernel Version = 5.15.0-1042-nvidia: OK
Linux perf_event_open syscall available: Fail
Sampling trigger event available: Fail
Intel(c) Last Branch Record support: Not Available
CPU Profiling Environment (process-tree): Fail
CPU Profiling Environment (system-wide): Fail

See the product documentation at https://docs.nvidia.com/nsight-systems for more information,
including information on how to set the Linux Kernel Paranoid Level.

What to look at: Fail on the CPU profiling lines means nsys will trace CUDA but cannot do CPU sampling. The usual cause is a paranoid level above 2 or a container without the needed capabilities; process-tree should be OK after lowering perf_event_paranoid.

cat /proc/sys/kernel/perf_event_paranoidKernel perf restriction level

What to look at: Ubuntu defaults to 4, which blocks unprivileged CPU sampling. nsys needs 2 or lower for process-tree CPU sampling (set with sysctl kernel.perf_event_paranoid=2), and system-wide sampling needs 1 or lower or root.

dcgmi profile --pauseFree GPU counters held by DCGM

Successfully paused profiling.

What to look at: DCGM profiling metrics hold the same hardware counters that ncu and nsys --gpu-metrics-devices need, so profiling fails or reports counters as busy until you pause. Remember dcgmi profile --resume afterwards, or monitoring dashboards show gaps in DCGM_FI_PROF_* metrics.

18. Scenarios

19. Searchable reference

Item	Meaning

20. Flashcards (click to flip)

21. Exam traps

Nsight Systems (nsys) = whole-application timeline. Nsight Compute (ncu) = one kernel in depth.
Work top-down: nsys first to find where time goes, then ncu on the dominant kernels.
ncu replays kernels and locks clocks: big overhead; not for whole production runs.
nsys overhead is low; keep the window short with --delay / --duration or a capture range.
Multi-node: srun nsys profile -o name_%q{SLURM_PROCID} app, not nsys profile srun app.
GPU gaps that line up with CPU file reads = input-bound, not a GPU or network problem.
A long NCCL kernel on most ranks can mean they're waiting for a straggler; compare ranks.
ERR_NVGPUCTRPERM = counter permission: run as admin or set NVreg_RestrictProfilingToAdminUsers=0.
Counter conflicts: pause DCGM profiling (dcgmi profile --pause) before ncu or nsys GPU metrics, then resume.
Nsight is for one job's performance; DCGM is for fleet monitoring and health; nccl-tests is for fabric bandwidth.
Reports: .nsys-rep (Systems) and .ncu-rep (Compute); record with the CLI on the node, view in the GUI elsewhere.
NVTX ranges label code regions on the timeline; NCCL emits them too.

22. Practice quiz (24 questions)

Score: 0 / 0