Lab 1.5 — Occupancy & the Roofline: Nsight Compute

← Course 4 syllabus · Module 1 · Prev: « Lab 1.4 · Next: Lab 2.1 »

Goal

The module’s capstone: learn to read a kernel the way a GPU engineer does, using the kernels you already own. Two instruments organize the reading. Occupancy — how many warps an SM can keep resident, and which resource (registers per thread, shared memory per block, block-size granularity) is the limiter — explains whether the machine can hide latency. The roofline model — attainable performance as \(P = \min(P_{\text{peak}},\; I \cdot BW)\), plotted against arithmetic intensity \(I\) — explains what a kernel is entitled to, and therefore whether optimizing compute or optimizing memory can possibly help. With both in hand, the Module 1 kernels (SAXPY, naive and tiled matmul, the reduction ladder) get placed on one chart, and the module closes with a disciplined optimization loop: hypothesis → counter → change → re-measure. Alongside time there is a third axis the big-iron literature underweights: performance per watt, read from nvidia-smi, treated here as a first-class metric even on a desktop part — as it will be for every frame budget in Modules 5–7.

Prerequisites

  • Labs 1.1–1.3: the kernels being analyzed, their ptxas -v outputs, and their saved Nsight Compute reports — this lab re-reads those artifacts as much as it takes new ones.
  • Lab 1.4’s nvidia-smi logging habit.
  • The Lab 1.1 device-properties table — every ceiling below comes from it or from a measured Module 1 bandwidth, never from a spec sheet.

Project & environment setup

  • No new kernels required; a small occupancy_report C++ target that calls cudaOccupancyMaxActiveBlocksPerMultiprocessor for each Module 1 kernel across block sizes is the one new program.
  • Rebuild variants come from flags, not code changes: --ptxas-options=-v (always), and -maxrregcount=<K> for the register-pressure experiment.
  • Nsight Compute drives the lab; the roofline needs the full section set:
ncu --set full -o labs/lab-1-5/captures/matmul_tiled ./matmul --tiled
  • Python: the same analysis applies to Numba kernels — ncu attaches to uv run python ... runs exactly as it did in Lab 1.1; one Numba kernel is included in the roofline chart to prove the point.

Where results go:

Artifact Path
Notes, roofline chart, occupancy tables, perf-per-watt table, optimization-loop log labs/lab-1-5/notes.md
Nsight Compute reports (one per kernel/variant) labs/lab-1-5/captures/
Block-size and -maxrregcount sweep CSVs, nvidia-smi logs labs/lab-1-5/benchmarks/

Background

  • Occupancy is a means, not a score. An SM has fixed budgets of registers, shared memory, and resident-warp slots; a kernel’s per-thread register count and per-block shared memory determine how many blocks fit, hence how many warps the scheduler can rotate through to hide latency. Achieved occupancy below theoretical points at launch geometry or workload imbalance. But a memory-bound kernel at 50% occupancy that already saturates the bus gains nothing from more warps — chasing occupancy for its own sake is the classic junior-GPU-engineer mistake this lab is designed to kill.
  • The roofline. With machine peak \(P_{\text{peak}}\) (FLOP/s) and memory bandwidth \(BW\) (B/s) — both taken from the device query and Module 1’s own measurements — a kernel of arithmetic intensity \(I\) (FLOP/byte) can at best achieve \(P = \min(P_{\text{peak}},\; I \cdot BW)\). The ridge point \(I^* = P_{\text{peak}}/BW\) splits the chart: left of it, memory-bound; right, compute-bound. Module 1’s kernels have hand-derivable intensities: SAXPY at \(I = 2/12 = 1/6\) FLOP/byte, the sum reduction at \(I = 1/4\) (one add per 4-byte load) — both far left. Naive matmul sits near \(I \approx 0.25\); tiled matmul rises to \(I \approx T/4\) for tile width \(T\) (Lab 1.2’s derivation), which is the roofline telling you why tiling worked: it moves the kernel rightward under the sloped roof.
  • Register pressure is the usual occupancy tax. ptxas -v reports registers per thread; capping with -maxrregcount (or __launch_bounds__) raises occupancy at the cost of register spills to local memory — a genuine trade to be measured, not a free win.
  • Perf per watt. Two kernels with equal runtime are not equal if one holds the GPU at higher power; in any thermally- or power-budgeted deployment the budget is joules, not just milliseconds — and a desktop 4090 makes the ranking easy to measure. Metric here: work per joule, from event-timed runs against time-aligned nvidia-smi power readings (nvidia-smi dmon or --query-gpu=power.draw) — coarse, but honest about the direction.

Tasks

  1. Static occupancy table. For each Module 1 kernel: registers/thread and shared mem/block from ptxas -v, then theoretical occupancy vs. block size from the occupancy API, tabulated. For each kernel, name the limiting resource in one sentence.
  2. Block-size sweep. Run SAXPY, tiled matmul, and the best reduction rung at block sizes 64 → 1024; event-time each. Plot runtime vs. block size next to theoretical occupancy vs. block size and mark where the curves stop tracking each other — that divergence is the “occupancy is a means” lesson.
  3. Register-pressure experiment. Rebuild the tiled matmul (your most register-hungry kernel) at 2–3 -maxrregcount values below its natural count. From ptxas -v: registers and spill bytes; from Nsight Compute: occupancy and local-memory traffic; from events: runtime. Find where the trade stops paying.
  4. Place the kernels on the roofline. Compute each kernel’s \(I\) by hand (show the arithmetic), take the ceilings from your device query and Module 1’s best measured bandwidth, and draw the chart (matplotlib is fine). Then open each kernel’s ncu --set full report and compare against Nsight Compute’s own roofline section — reconcile your placement with the tool’s, including any definitional differences in how bytes are counted (DRAM vs. L2-level intensity).
  5. The optimization loop, run once for real. Pick the Module 1 kernel with the largest gap to its roof. Write the hypothesis first (“bounded by X because counter Y says Z”), make one change, re-profile, and log the loop in notes.md — including if the change lost, which is a finding, not a failure.
  6. Perf-per-watt table. For SAXPY, tiled matmul, and the best reduction: sustained run under an nvidia-smi dmon log, average power over the run window, compute work-per-joule. One paragraph: which kernel is most efficient, and does the ranking match the time-only ranking?

Deliverable & expected results

  • The roofline chart with all Module 1 kernels placed (hand-computed and tool-reported), the occupancy and sweep tables, the logged optimization loop, and the perf-per-watt table.
Quantity Predicted Measured
SAXPY roofline position far left — \(I = 1/6\) FLOP/byte, memory-bound with no escape; attainable \(P \approx I \cdot BW\) from your measured bandwidth …
Reduction (best rung) position \(I = 1/4\), same left region — its achieved bandwidth should nearly touch the sloped roof …
Tiled vs. naive matmul movement rightward shift ∝ tile width (\(I \approx T/4\) vs. \(\approx 1/4\)) — read straight off the chart …
Block-size sweep shape flat plateau once occupancy suffices to hide latency; penalties only at the small-block edge (and possibly the largest sizes via granularity effects) …
-maxrregcount trade occupancy up, but spills eventually swamp the gain — a non-monotone curve with an interior best …
Perf-per-watt ranking plausibly not identical to the time ranking — direction to verify, not assume …

Profiling & performance

Nsight Compute is the subject, not just the tool. For each kernel keep one --set full report and learn three of its pages cold: Speed Of Light (the summary verdict + roofline), Occupancy (theoretical vs. achieved, limiter attribution), and Scheduler/Warp State statistics (what warps were waiting on — the stall reasons that justify the word “latency-bound” when neither roof is hit). Cross-check one kernel end-to-end: does the SOL memory percentage agree with your event-timed effective-bandwidth number within reason? If not, find the definitional gap before trusting either.

Analysis & reconciliation

The module’s closing essay, one page: for each Module 1 kernel — position on the roofline (hand vs. tool), the limiting resource, achieved vs. theoretical occupancy and whether the gap mattered, and the one change that would (or provably wouldn’t) move it. Reconcile the block-size plateau against the occupancy table, the spill trade against local-memory counters, and the perf-per-watt ranking against the power logs. End with the transferable statement of the loop — counter, hypothesis, single change, re-measure — because that loop, not any CUDA syntax, is what Modules 5–7 reuse.

Going further

  • Add an __launch_bounds__ annotation matching your best sweep result and confirm ptxas holds registers under the implied cap — the production spelling of today’s experiment.
  • Roofline a cupy.matmul (cuBLAS) run from its Nsight Compute report and see how close to the ridge a professional kernel sits.
  • Hold this lab’s method in mind for Module 6: the same analysis style — ceilings, limiters, one-change loops, watts — returns for graphics workloads, where Nsight Compute’s roofline is replaced by GPU captures and counter sets but the reasoning is identical (Lab 6.4 is this lab wearing a render pass).