<<<CUDA C++ Coursegrid · block · warp · lane
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.

ToolWhat it showsWhen to use it
Nsight Systems (nsys)Whole-app timeline: kernels, copies, CPU, API callsStep one. Find who is slow, whether there are gaps, whether overlap is actually happening
Nsight Compute (ncu)Hardware-counter detail for a single kernelStep two. You know which kernel is the bottleneck and need to dig inside it
compute-sanitizerOut-of-bounds, uninitialized memory, races, leaksWhen correctness is in doubt, or results are unstable
nvidia-smi dmonUtilization, power, temperature, clocksWhen you suspect throttling or multi-process contention

Step one: nsys for the whole app

Capture a timeline
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

Commands you will use
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
MetricMeaningHow to read it
Compute ThroughputSM compute utilization> 60% means compute-bound: optimize instructions and the algorithm
Memory ThroughputDevice-memory bandwidth utilization> 70% means bandwidth is nearly saturated: the only lever is less data movement
Achieved OccupancyFraction of warps actually residentFar below theoretical → tail effect or load imbalance
sectors / requestsAverage sectors per requestIdeal is 4; near 32 means the loads are completely uncoalesced
Bank ConflictsShared-memory conflict countAny nonzero is worth trying to pad away
Warp Stall ReasonsWhat warps are waiting onSee the next table. This is the most informative metric
Stall reasonMeaningWhat to do
Long ScoreboardWaiting on global memoryRaise occupancy, improve coalescing, add ILP
Short ScoreboardWaiting on shared memoryEliminate bank conflicts
BarrierStuck at __syncthreads()Fewer barriers, more even work inside the block
MIO ThrottleMemory-instruction queue is fullVectorize loads to cut the instruction count
Math Pipe ThrottleCompute pipes are queuedGood sign: you are actually computing
Not SelectedAnother warp was scheduled firstGood 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 algorithm
Roofline model
Generate a roofline analysis
1ncu --set roofline -o roofline_report ./app2ncu-ui roofline_report.ncu-rep      # the GUI plots where each kernel lands

compute-sanitizer: valgrind for CUDA

Four checking modes
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.cu

A reusable optimization loop

  1. 1.Build a baseline: measure stable time and effective bandwidth with cudaEvent, and put that in a script.
  2. 2.Compute the theoretical ceiling: bytes / device-memory bandwidth = shortest possible time. The gap is then obvious.
  3. 3.Locate with nsys: which kernel dominates? Are there gaps and serialization?
  4. 4.Diagnose with ncu: classify from SOL, pinpoint from stall reasons, check coalescing with sectors/requests.
  5. 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. 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.