Profiling & Debug — Nsight Systems, Nsight Compute, NVTX, CUPTI

Topics We'll Cover

A practical walkthrough of NVIDIA's profiling stack. Each tool answers a different question; the skill is knowing which one to reach for first.

• The Profiling Hierarchy — Fleet, System, Kernel, SASS
• NVTX — Annotate First, Profile Second
• Nsight Systems — The System Timeline
• Reading a Timeline — Bad vs Good
• Nsight Compute — Kernel Deep Dive
• Roofline Analysis — Memory- vs Compute-Bound
• Common Stalls and Fixes
• Distributed Profiling — Per-Rank, NCCL
• CUPTI — Programmatic Profiling
• Fleet & Operational Tools — DCGM, nvbandwidth, nvitop
• Interactive: Pick Your Tool

The Profiling Hierarchy

Profiling is not one tool, it's a stack. A different question is asked at every level — and a different tool answers it. Picking the wrong tool wastes hours: nobody finds a bank conflict in nvidia-smi, and nobody finds a flaky power supply in Nsight Compute.

The four levels in plain English

• Fleet — "Is the cluster healthy?" DCGM, dcgmi, dcgm-exporter feeding Prometheus. XID errors, ECC counts, throttling reasons, power, link state. Always-on; alerts fire long before users complain.
• System — "Where is one rank's wall-clock going?" Nsight Systems renders a unified timeline of CPU threads, CUDA streams, kernels, NCCL collectives, OS scheduling, and your NVTX ranges. Best at finding gaps, stalls, oversubscription.
• Kernel — "Is this specific kernel doing what I think?" Nsight Compute (ncu) measures occupancy, achieved bandwidth, achieved FLOPS, warp stall reasons, register and shared-memory usage, source-correlated SASS.
• SASS / source — "Why does the compiler emit those instructions?" cuobjdump and nvdisasm on the cubin, plus --ptxas-options=-v at compile. The bedrock when nothing higher up explains the behaviour.

NVTX — Annotate First, Profile Second

NVTX (NVIDIA Tools Extension) is the bridge between your code's mental model and what the profiler sees. It is just two things: range markers (push/pop) and category labels. The cost at runtime is essentially zero when no profiler is attached — ship NVTX in production code without guilt.

import torch
import torch.cuda.nvtx as nvtx

def train_step(model, batch, optimizer, criterion):
    nvtx.range_push("step")

    nvtx.range_push("forward")
    logits = model(batch["input_ids"])
    loss = criterion(logits, batch["labels"])
    nvtx.range_pop()                         # /forward

    nvtx.range_push("backward")
    loss.backward()
    nvtx.range_pop()                         # /backward

    nvtx.range_push("optim")
    optimizer.step()
    optimizer.zero_grad(set_to_none=True)
    nvtx.range_pop()                         # /optim

    nvtx.range_pop()                         # /step
    return loss.item()

#include <nvtx3/nvToolsExt.h>

void run_block(int layer_id) {
    char name[64];
    snprintf(name, sizeof(name), "block.%d", layer_id);
    nvtxRangePushA(name);

    nvtxRangePushA("attn");
    launch_attention_kernel(...);
    nvtxRangePop();                          // /attn

    nvtxRangePushA("mlp");
    launch_mlp_kernel(...);
    nvtxRangePop();                          // /mlp

    nvtxRangePop();                          // /block.N
}

Nsight Systems — System Timeline

Nsight Systems (nsys) is the system-wide profiler. One capture covers every CPU thread, every CUDA stream, every kernel launch, every NCCL collective, OS scheduling, and your NVTX ranges — correlated on a single time axis. It is the first tool to reach for when something is "slow".

# Single rank, full coverage:
nsys profile \
    -t cuda,nvtx,osrt,cudnn,cublas,nccl \
    --capture-range=cudaProfilerApi  \
    --capture-range-end=stop          \
    -o train_run                      \
    python train.py

# Then open in Nsight Systems UI:
nsys-ui train_run.nsys-rep

# Or generate a CLI summary without the UI:
nsys stats train_run.nsys-rep | head -60

The trace categories

• cuda — every CUDA Runtime & Driver API call: launches, copies, synchronisations, stream and event ops. Mandatory.
• nvtx — user and library NVTX ranges. Pair with the cuda rows on the timeline to read intent.
• osrt — OS runtime: file IO, mutexes, sleep, futex. Reveals dataloader stalls and Python GIL contention.
• cudnn / cublas — library-level NVTX from the maths kernels themselves.
• nccl — collective ops with size and rank metadata. Essential for distributed.

What you actually look for

Reading a Timeline

What does a "good" timeline look like, vs a "bad" one? The difference is rarely subtle once you know what to look for.

The features to read

• Stream rows — each CUDA stream gets its own row. Multiple busy rows in parallel = healthy overlap; one row of work + others empty = serialisation.
• NVTX bands — coloured strips above the streams; named after your range pushes. Use them as section headers.
• Kernel duration — click any GPU bar; the inspector shows kernel name, grid/block, registers, shared memory, and which API call launched it.
• Occupancy plot — the GPU SM activity heatmap below the streams. Ideally a flat green band; gaps mean idle SMs.

Nsight Compute — Kernel Deep Dive

Nsight Compute (ncu) is the kernel-level microscope. Where Nsight Systems gives you the timeline, ncu gives you the inside of one bar on that timeline: which exact PerfMon counters fired, how many warps stalled and on what, what fraction of theoretical bandwidth and FLOPS you actually achieved.

# Profile every kernel with the full metric set, write report to file:
ncu --set full -o kernel_full python repro.py

# Most kernels you don't care about. Filter by name:
ncu --set full -k "matmul|attention" -o focused python repro.py

# Only the first occurrence of each kernel:
ncu --set full --launch-skip 0 --launch-count 1 -o once python repro.py

# Source-correlated SASS view (huge for finding the line that stalls):
ncu --set full --import-source yes -o with_src python repro.py

# Then open in the UI:
ncu-ui kernel_full.ncu-rep

What ncu reports per kernel

Roofline Analysis

The roofline plot is the single best mental model for "is this kernel slow because of memory or because of compute?". It pins the kernel on a chart with two ceilings: the memory bandwidth ceiling (left, sloped) and the compute ceiling (right, flat). Whichever is closer wins; you cannot go above either.

How to read it

• X-axis = arithmetic intensity (FLOPs per byte read from HBM). Property of the algorithm and tile size, not the GPU.
• Y-axis = achieved throughput in FLOPS. Property of this run.
• Sloped roof = memory bandwidth limit; flat roof = compute (tensor-core) limit.
• Ridge point = peak_FLOPS / peak_bandwidth. On H100 SXM5 with dense FP8 (~1979 TFLOPS / 3.35 TB/s HBM3), the ridge sits around AI ≈ 590 FLOP/byte (roughly double that with structured sparsity).

Practical use — what the dot tells you

Common Stalls and Fixes

The Warp State section of ncu reports the average reason a warp was not eligible to issue. Each reason maps to a class of fix. Knowing the table by heart is what separates "the kernel is slow" from "the kernel needs cp.async with a deeper pipeline".

Distributed Profiling

One rank tells you about that rank. Distributed problems — stragglers, ring imbalance, cross-node latency — only show up when you correlate timelines across ranks. Nsight Systems has first-class support for this.

# torchrun: each rank gets its own .nsys-rep, named by RANK env var:
nsys profile -t cuda,nvtx,nccl -o "out_rank%q{RANK}" \
    torchrun --nproc_per_node=8 --nnodes=2 train.py

# SLURM srun: same idea with SLURM_PROCID:
srun --ntasks=16 nsys profile -t cuda,nvtx,nccl \
    -o "out_%q{SLURM_PROCID}" python train.py

# Diagnostic env vars to capture alongside the trace:
export NCCL_DEBUG=INFO              # protocol, ring, algo, channel choice
export NCCL_DEBUG_SUBSYS=COLL,ENV   # filter what NCCL prints
export NCCL_TOPO_DUMP_FILE=topo.xml # NCCL's view of the fabric — gold for "why slow ring"

Multi-report timeline

Nsight Systems' UI lets you load several .nsys-rep files together: File → Open Multiple. The tool aligns them on a common time axis (using NCCL collective endpoints as anchors). What you can then see at a glance:

• Stragglers — one rank's all-reduce bar is consistently longer than its peers; that rank started the collective late, the others were waiting.
• Imbalanced rings — collective duration scales with the slowest pair, not the average. Cross-node hops show up as longer NCCL bars.
• PCIe-bound nodes — nodes without NVSwitch show fatter all-reduce bars even when GPUs are idle on either side.
• Dataloader skew — ranks reading from a slow filesystem fall behind; visible as one rank's compute bars shifted right.

CUPTI — Programmatic Profiling

CUPTI (CUDA Profiling Tools Interface) is the C API behind every NVIDIA profiler. It surfaces every CUDA event — kernel launches, memcpys, sync ops, stream activity — plus the GPU PerfMon counters used by Nsight Compute. You won't write CUPTI directly often; you almost certainly already use it through torch.profiler and friends.

#include <cupti.h>
#include <stdio.h>

static void CUPTIAPI on_buffer_request(uint8_t **buf, size_t *sz, size_t *maxn) {
    *sz = 8 * 1024 * 1024;
    *buf = (uint8_t*)aligned_alloc(8, *sz);
    *maxn = 0;
}

static void CUPTIAPI on_buffer_complete(CUcontext ctx, uint32_t streamId,
                                          uint8_t *buf, size_t sz, size_t valid) {
    CUpti_Activity *rec = NULL;
    while (cuptiActivityGetNextRecord(buf, valid, &rec) == CUPTI_SUCCESS) {
        if (rec->kind == CUPTI_ACTIVITY_KIND_KERNEL ||
            rec->kind == CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL) {
            CUpti_ActivityKernel9 *k = (CUpti_ActivityKernel9*)rec;
            printf("%-40s  %.3f ms  grid=(%u,%u,%u) regs=%u\n",
                   k->name,
                   (k->end - k->start) / 1.0e6,
                   k->gridX, k->gridY, k->gridZ,
                   k->registersPerThread);
        }
    }
    free(buf);
}

int main(void) {
    cuptiActivityRegisterCallbacks(on_buffer_request, on_buffer_complete);
    cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL);
    cuptiActivityEnable(CUPTI_ACTIVITY_KIND_MEMCPY);

    /* ... run your CUDA program here ... */

    cuptiActivityFlushAll(1);
    return 0;
}

Fleet & Operational Tools

Nsight is for development. In production you don't want to attach a profiler — you want continuous, low-overhead telemetry that alerts when something drifts. That's where DCGM, nvbandwidth, nvitop, and the stress testers live.

What DCGM surfaces that nvidia-smi doesn't

• XID error codes — numbered driver-detected GPU faults (e.g. XID 79 = "fell off the bus"). Hard signal of hardware failure.
• ECC error counts — corrected (SBE) and uncorrected (DBE), per memory partition. Trends matter more than absolute values.
• Throttling reasons — thermal, power, SW slowdown, HW slowdown — with millisecond-resolution history.
• Health snapshots — aggregate "is this GPU OK?" boolean per second.
• NVLink errors — replays, recoveries, lane-down events.

# On every GPU node: dcgm-exporter as a sidecar / DaemonSet
docker run -d --name dcgm-exporter --gpus all --cap-add SYS_ADMIN \
    -p 9400:9400 nvcr.io/nvidia/k8s/dcgm-exporter:latest

# Prometheus scrapes :9400/metrics every 15s
# Grafana dashboard 12239 (DCGM Exporter) is a sane default

# Pre-flight: full diagnostic before reintroducing a node
sudo dcgmi diag -r 3          # level 3 = ~5 min, full battery
nvbandwidth -t host_to_device_memcpy_ce
nvbandwidth -t device_to_device_memcpy_read_ce

Interactive: Pick Your Tool

Pick the question you actually have. The planner picks the tool, the exact command, what to look for, and the next tool to reach for if the first one doesn't surface the answer.