Stage 5 · Production and AI13 / 15about 24 min
Profiling: replace guesses with data
How Nsight Compute and Nsight Systems split the work, the metrics that matter, and how to read a roofline
After this lesson you can
- Know which questions nsys answers and which ncu answers
- Memorize the five metrics to check on every kernel
- Read a roofline plot and pick an optimization direction from it
- Use compute-sanitizer to find out-of-bounds access and races
The most common CUDA optimization failure is speeding up something that is not the bottleneck. You spend two days cutting a kernel's instruction count by 30%, and end-to-end time does not change, because the kernel was already bound by device-memory bandwidth. There is one rule: measure, change, measure again. NVIDIA's tools split the work cleanly. Using the right one saves a lot of time.
| Tool | What it shows | When to use it |
|---|---|---|
Nsight Systems (nsys) | Whole-app timeline: kernels, copies, CPU, API calls | Step one. Find who is slow, whether there are gaps, whether overlap is actually happening |
Nsight Compute (ncu) | Hardware-counter detail for a single kernel | Step two. You know which kernel is the bottleneck and need to dig inside it |
compute-sanitizer | Out-of-bounds, uninitialized memory, races, leaks | When correctness is in doubt, or results are unstable |
nvidia-smi dmon | Utilization, power, temperature, clocks | When you suspect throttling or multi-process contention |
Step one: nsys for the whole app
shell
1nsys profile -o report --stats=true --force-overwrite=true ./app2 3# Print a stats summary from the command line4nsys stats report.nsys-rep5 6# Open the timeline in the GUI (this is where you see overlap across streams)7nsys-ui report.nsys-rep- ▸Large gaps between kernels → the CPU side is the bottleneck. Consider CUDA Graph or fewer synchronizations.
- ▸Copies and compute never overlap → check whether you are on the default stream, or host memory is not pinned.
- ▸One kernel takes 80% of the time → that is the target. Go to step two.
- ▸Many tiny kernels → consider operator fusion.
Step two: ncu for the details
shell
1# Full analysis (slow: the kernel is replayed many times)2ncu --set full -o profile ./app3 4# Profile only the first 3 launches of a named kernel5ncu --kernel-name matmulTiled --launch-count 3 ./app6 7# Pull only the metrics you care about; much faster8ncu --metrics \9 sm__throughput.avg.pct_of_peak_sustained_elapsed,\10 gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed,\11 sm__warps_active.avg.pct_of_peak_sustained_active,\12 l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\13 l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum \14 ./app15 16# Let the tool print optimization advice17ncu --set full --section SpeedOfLight --section Occupancy ./app| Metric | Meaning | How to read it |
|---|---|---|
| Compute Throughput | SM compute utilization | > 60% means compute-bound: optimize instructions and the algorithm |
| Memory Throughput | Device-memory bandwidth utilization | > 70% means bandwidth is nearly saturated: the only lever is less data movement |
| Achieved Occupancy | Fraction of warps actually resident | Far below theoretical → tail effect or load imbalance |
| sectors / requests | Average sectors per request | Ideal is 4; near 32 means the loads are completely uncoalesced |
| Bank Conflicts | Shared-memory conflict count | Any nonzero is worth trying to pad away |
| Warp Stall Reasons | What warps are waiting on | See the next table. This is the most informative metric |
| Stall reason | Meaning | What to do |
|---|---|---|
Long Scoreboard | Waiting on global memory | Raise occupancy, improve coalescing, add ILP |
Short Scoreboard | Waiting on shared memory | Eliminate bank conflicts |
Barrier | Stuck at __syncthreads() | Fewer barriers, more even work inside the block |
MIO Throttle | Memory-instruction queue is full | Vectorize loads to cut the instruction count |
Math Pipe Throttle | Compute pipes are queued | Good sign: you are actually computing |
Not Selected | Another warp was scheduled first | Good sign: there is enough parallelism |
Roofline: put the kernel on a chart
Performance (GFLOP/s)
▲
19500┤ ┌────────────────── peak compute (compute roof)
│ ╱
│ ╱ ← slope region: bandwidth-bound
│ ╱ performance = bandwidth × arithmetic intensity
│ ╱
│ ● ╱ ● tiled matmul (still room to climb)
│ ╱
│ ● ╱ ● vector add (on the slope = bandwidth saturated)
└──────┴──────────────────────────▶ arithmetic intensity (FLOP/Byte)
12.5
ridge point = peak compute / peak bandwidth
on the slope → bandwidth is saturated; raise intensity (fusion, tiling, reuse)
below the slope → room to optimize: uncoalesced access / low occupancy / conflicts
on the plateau → compute is saturated; consider Tensor Core or a different algorithmshell
1ncu --set roofline -o roofline_report ./app2ncu-ui roofline_report.ncu-rep # the GUI plots where each kernel landscompute-sanitizer: valgrind for CUDA
shell
1# Out-of-bounds and illegal addresses (the one you use most)2compute-sanitizer --tool memcheck ./app3 4# Races in shared or global memory5compute-sanitizer --tool racecheck ./app6 7# Reads of uninitialized memory8compute-sanitizer --tool initcheck ./app9 10# Device-memory leaks11compute-sanitizer --tool synccheck ./app12 13# To map errors to source lines, compile with -lineinfo (lighter than -G, does not turn off optimization)14nvcc -O3 -lineinfo -arch=sm_80 -o app main.cuA reusable optimization loop
- 1.Build a baseline: measure stable time and effective bandwidth with cudaEvent, and put that in a script.
- 2.Compute the theoretical ceiling: bytes / device-memory bandwidth = shortest possible time. The gap is then obvious.
- 3.Locate with nsys: which kernel dominates? Are there gaps and serialization?
- 4.Diagnose with ncu: classify from SOL, pinpoint from stall reasons, check coalescing with sectors/requests.
- 5.Change one thing: one variable per round, remeasure immediately. Change three things at once and you never know which one helped (or cancelled the others).
- 6.Check correctness: after every round, compare against a reference and run compute-sanitizer. A fast wrong kernel is worthless.
Self-check
Answer in your head first, then open the solution. If you cannot, reread that section.