GPU Architecture · Execution & Operations
What metadata travels with your kernel, and exactly which data source each tool reads — so a symptom points you at the right instrument.
This is where the arc pays off. You know how a launch runs (L7), how data, completion, and errors flow (L8), and how the GPU is shared (L9). When something is wrong or slow, you reach for a tool — but the tools aren't interchangeable, because they watch the GPU through different windows. Some read cheap driver counters, some trace events on a timeline, some read expensive hardware performance counters, some instrument your binary. Knowing which data source a tool uses is what tells you when to reach for it.
Match the question to the instrument: "is it even using the GPU?" is a driver-counter question (nvidia-smi); "where does wall-clock go?" is a tracing question (Nsight Systems); "why is this kernel slow?" is a hardware-counter question (Nsight Compute); "is it wrong?" is an instrumentation question (compute-sanitizer). Different windows, different data.
A tool can only show you source lines if you compiled the mapping in. Two flags, and the difference matters:1
-lineinfo — for profiling.debug_line) — without disabling optimization.-G — for debuggingcuda-gdb to inspect variables and single-step.Mechanically: a profiler samples a program-counter (a SASS instruction address), looks it up in the cubin's .debug_line table to recover file+line, and aggregates its counters per source line. No line info compiled in → the source column is blank and you're reading raw SASS and mangled names. A release build strips these sections, which is why a stripped binary profiles as addresses only.1
The three performance tools form a ladder — each narrower in scope, deeper in detail, and costlier to run than the last. You walk down it: start cheap and wide, descend only as far as the question demands.
Is it used? Memory, power, temp, clocks, throttling, which processes.
Where does wall-clock go? Gaps, poor overlap, CPU-bound stretches, tiny kernels.
Why is this kernel slow? Occupancy, memory vs compute bound, warp stalls.
What separates them is entirely their data source:
| Tool | What it shows | Data source | Cost |
|---|---|---|---|
| nvidia-smi / NVML |
Utilization, memory, power, clocks, temp, throttle reasons, processes. | Driver counters & on-chip sensors | ~free, live |
| Nsight Systems |
Unified CPU+GPU timeline: API calls, kernels, memcpies, overlap, gaps; a launch linked to its kernel. | Tracing via CUPTI (event records + timestamps + correlation IDs) | low |
| Nsight Compute |
Per-kernel depth: achieved occupancy, memory throughput vs peak (roofline), warp stall reasons, cache/sector efficiency. | Hardware performance counters + PC/warp-state sampling | high |
Three details worth carrying:
nvidia-smi "GPU-Util" trap. That percentage is the fraction of time at least one kernel was running — time-busy, not saturation. A kernel using a single SM of your twenty reports 100%. So nvidia-smi tells you the GPU is occupied, never that it's well used — for that you need occupancy from Nsight Compute.3cudaLaunchKernel to the exact GPU kernel it started — that's the line drawn between the CPU and GPU lanes on the timeline.4ncu re-runs your kernel multiple times — one "pass" per group of counters — saving and restoring its memory between passes so each replay sees identical inputs, and it flushes caches and serializes launches for determinism. That's why a profiled kernel runs far slower than normal, and why you point it at one kernel, not a whole app.5Nsight Compute's most illuminating output samples the warp program counter to report why warps were idle — the very stall reasons from Lesson 2. stall_long_scoreboard = waiting on global/local memory (the bandwidth wall of Lesson 4); stall_barrier = waiting at a __syncthreads() (Lesson 6); stall_short_scoreboard = shared-memory/MIO latency. This is where every physical lesson becomes a number you can read off your own kernel.5
Slow and wrong are different questions, and the counters above won't catch wrong. Two tools work by different means entirely:
compute-sanitizer — instruments your binary to catch correctness bugs: memcheck (out-of-bounds / misaligned access), racecheck (shared-memory data races — you met this in Lesson 6), initcheck (reads of uninitialized global memory), synccheck (illegal __syncthreads() use). It's the tool that finds the bug a data race hides, deterministically, where a lucky run passes.6cuda-gdb — a debugger back-end that halts device execution so you can set breakpoints in a kernel, switch focus to a specific block / warp / lane, and inspect its registers and memory. Needs a -G build for full source-level debugging. Reads state; doesn't sample counters.7All of these run in your WSL2 setup. Two practical notes: build with -lineinfo whenever you profile so Nsight can point at source lines; and Nsight Compute's hardware counters require performance-counter access — if ncu reports ERR_NVGPUCTRPERM, that's a permissions gate to enable, not a broken install. compute-sanitizer and nvidia-smi need no special setup.
The whole lesson, compressed to "symptom → reach for → because of the data it reads":
| Symptom | Reach for | Because it reads… |
|---|---|---|
| Is the GPU even used? OOM? throttling? | nvidia-smi | driver counters & sensors |
| GPU shows 100% but it's slow | → Nsight Compute | real occupancy (HW counters), not time-busy |
| Where's wall-clock going? bad overlap, tiny kernels, CPU-bound? | Nsight Systems | CUPTI trace + correlation IDs |
| This one kernel is slow — memory- or compute-bound? | Nsight Compute | HW counters (roofline) |
| Why are warps idle? | Nsight Compute | PC / warp-state sampling (stall reasons) |
| Wrong results / crash — bad pointer, race, uninit, bad sync? | compute-sanitizer | binary instrumentation |
| Step through device code, inspect a lane's registers | cuda-gdb (-G) | debugger back-end + DWARF |
| Crash surfaced on an innocent later call (L8 sticky error) | compute-sanitizer | instrumentation pins the real kernel + address |
That's the whole decision tree. For the full ecosystem — every tool in the course plus the ones we didn't reach (nsys/ncu siblings, CUPTI, NVTX, DCGM, the monitors, and current-vs-deprecated status), each with its data source and best article — see the companion CUDA Tools & Further Reading card.
From memory — one click locks each answer.
Data sourcenvidia-smi reports your GPU at 100% utilization, yet the program is slow. This is because GPU-Util —
GPU-Util is the fraction of time at least one kernel was running — time-busy, not capacity. A kernel occupying a single SM of many still shows 100%. For true utilization you need achieved occupancy from Nsight Compute's hardware counters.
Tool choiceYour app spends most of its wall-clock time somewhere, but no single kernel looks slow. You reach for —
"Where does wall-clock go" is a tracing question: Nsight Systems (CUPTI trace) shows the CPU+GPU timeline — gaps, missing overlap, CPU-bound stretches, tiny back-to-back kernels. Nsight Compute is for when you've already found the one slow kernel.
MechanismNsight Compute re-runs your kernel several times because —
Hardware limits how many performance counters collect at once, so ncu replays the kernel once per group of counters ("kernel replay"), saving/restoring memory so each pass sees identical input. That replay — plus cache flushes and serialized launches — is why profiled kernels run much slower.
Nsight Compute — Kernel Profiling Guide (kernel replay, hardware counters, warp-state sampling, roofline): read it here. Then skim the CUPTI overview (Activity/Callback APIs, correlation IDs) to see what Nsight Systems is built on: docs.nvidia.com/cupti.
That closes the execution & operations arc — and the tools in this lesson are exactly the ones the optimization arc will put in your hands. Next, the payoff we've been building toward since Lesson 6: a real tiled matmul kernel — load a tile into shared memory once, reuse it many times — measured against your 320 GB/s, then profiled under Nsight Compute so you read stall_long_scoreboard dropping on your own code as the tiling takes hold. Every lesson so far becomes a number you can move. Say the word when you're ready to build it.
--generate-line-info (-lineinfo): "Generate line-number information for device code" (optimizations intact; for profiling). --device-debug (-G): generates device debug info and, without --dopt, "turns off all optimizations… not intended for profiling." CUDA Binary Utilities — cubins are ELF with DWARF .debug_* sections; .debug_line maps address → source line. nvcc guide · binary utilitiesnvmlUtilization_t.gpu is the "percent of time over the past sample period during which one or more kernels was executing." Time-busy, not saturation: a kernel on one SM reports 100%. (Use SM-activity/occupancy metrics for true utilization.) NVML APIcorrelationId that "matches the correlation ID of the associated kernel, memcpy, or memset activity record that the API call initiated"; kernel records carry GPU-timestamp start/end. This is how Nsight Systems links the host launch to the GPU kernel. CUPTI Activity APImemcheck (out-of-bounds/misaligned), racecheck (shared-memory races), initcheck (uninitialized global reads), synccheck (invalid synchronization); runtime binary instrumentation; replaces the deprecated cuda-memcheck. docs.nvidia.com/compute-sanitizercuda device sm warp lane block thread focus; info cuda warps/lanes; full source-level debugging needs a -G build. docs.nvidia.com/cuda/cuda-gdb