2. Build a trustworthy baseline

Fix the experiment

Record the dimensions that change code generation or library selection. For a GEMM these are $M,N,K$, layouts, dtype, epilogue (the post-accumulation step — bias, activation, scaling, store conversion), and alignment. For a model they include batch, sequence length, hidden sizes, precision policy, checkpointing, and graph/compile mode.

The primary rate is usually

$$\text{throughput}=\frac{\text{useful work}}{\text{wall time}},$$

where useful work may be tokens, images, requests, cells, or bytes. Define it so an optimization cannot “win” by silently doing less work.

Warm-up and timing

Warm-up must cover five effects, each of which makes the first iterations artificially slow for a distinct mechanical reason: lazy module loading (the CUDA driver does not load and link a kernel's binary until the first launch that needs it); JIT compilation (if the binary ships portable PTX instead of architecture-exact machine code, the driver compiles it to native instructions on first use, which is slow); autotuning (libraries like cuBLAS/cuDNN benchmark several internal kernel candidates on first call for a given shape and cache the winner); allocator growth (CUDA's caching memory allocator requests new device memory from the driver the first time a size is needed, which is far slower than reusing an already-grown pool); and cache creation (persistent kernels, workspaces, and lookup tables get built once and reused thereafter). Then synchronize around the measured region:

for (int i = 0; i < warmup; ++i) step();
cudaDeviceSynchronize();

auto t0 = std::chrono::steady_clock::now();
for (int i = 0; i < iterations; ++i) step();
cudaDeviceSynchronize();
auto t1 = std::chrono::steady_clock::now();

CUDA events measure elapsed device-stream time without turning every operation into a host synchronization:

cudaEventRecord(start, stream);
step(stream);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&milliseconds, start, stop);

Use wall time for end-to-end truth and events to decompose GPU work. They answer different questions. CUDA calls are asynchronous — the host thread returns as soon as work is enqueued, not when it finishes — so host-side wall-clock timing around an individual call mostly measures dispatch latency. CUDA events are timestamps recorded directly on the GPU stream by the hardware and driver, so they measure real device execution time without forcing a host synchronization for every call.

Report a distribution

A single number hides variance from thermal state, OS scheduling, and contention. Report enough of the distribution that a reader can judge whether a change is real:

Locking clocks can reduce noise, but it changes the operating regime. If clocks are not locked, record sampled clocks and compare utilization ratios more than single-run nanoseconds in replayed ncu captures.

Baseline checklist

Capture the environment alongside the numbers, since driver, clock state, and topology all shift results without changing the code:

nvidia-smi --query-gpu=name,driver_version,pstate,clocks.sm,clocks.mem,\
temperature.gpu,power.draw,memory.total --format=csv
ncu --version
nsys --version

Also record CPU model, NUMA topology, memory capacity, framework/toolkit versions, and whether the environment is native Linux, Windows, a container, or WSL. Host scheduling and profiler permissions differ materially across them.

Capacity is part of performance

Track peak allocated and reserved device memory. Host or unified-memory spill can make a configuration appear to “run” while reducing throughput by orders of magnitude. A valid baseline states whether the complete steady-state workload, including optimizer state and periodic passes, fits.