Performance Analysis

Assess → Parallelize → Optimize → Deploy (from the CUDA C++ Best Practices Guide).

21.1 The APOD cycle

Assess → Parallelize → Optimize → Deploy (from the CUDA C++ Best Practices Guide). Profile first; optimize the biggest bottleneck; repeat.

21.2 Nsight Systems — system-wide timeline

Shows CPU threads, CUDA API calls, kernel launches, memory copies, stream concurrency, NVTX ranges, gaps/idle time. Answers: *Is the GPU busy? Are copies overlapping? Where is CPU overhead?* NVTX annotations:

nsys profile -o report --stats=true ./app
nsys-ui report.nsys-rep

#include <nvtx3/nvToolsExt.h>
nvtxRangePushA("preprocess");
preprocess<<<...>>>();
nvtxRangePop();

21.3 Nsight Compute — kernel-level analysis

Metrics: SM throughput, memory throughput (DRAM/L2/L1/shared), achieved occupancy, warp stall reasons (memory dependency, execution dependency, barrier, not selected), instruction mix, bank conflicts, uncoalesced accesses, roofline chart, source-level hot spots (with -lineinfo).

ncu --set full -o kernel_report ./app
ncu --kernel-name matmulTiled --launch-skip 2 --launch-count 1 ./app
ncu-ui kernel_report.ncu-rep

21.4 Roofline model

Arithmetic intensity (AI) = FLOPs / bytes moved from DRAM.

Ridge point = peak FLOP/s ÷ peak bandwidth (e.g., H100 FP32 ≈ 67 TFLOP/s ÷ 3.35 TB/s ≈ 20 FLOP/byte).

Below the ridge → memory-bound (optimize data movement); above → compute-bound (optimize instructions, use Tensor Cores).

vecAdd: 1 FLOP per 12 bytes → AI ≈ 0.08 → heavily memory-bound. Large GEMM: AI grows with N → compute-bound.

Performance (FLOP/s)
   │            ______________________  peak compute
   │          /
   │        /   ← memory-bound region (slope = bandwidth)
   │      /
   └──────────────────────────────────── Arithmetic intensity (FLOP/byte)

21.5 Bottleneck categories

high DRAM throughput, low SM utilization coalescing, caching in shared memory, fusion, lower precision, compression low throughput on both, stalls on memory more parallelism (occupancy/ILP), prefetching, async copies high SM/FMA pipe utilization algorithmic improvements, Tensor Cores, fast math intrinsics divergence, many integer/index ops restructure branches, simplify indexing, unroll Launch/CPU overhead many tiny kernels, gaps in timeline kernel fusion, CUDA graphs, persistent kernels copies dominate timeline keep data on GPU, overlap, pinned memory