Correct GPU Profiling and Benchmarking

A patch arrives with one exciting line: triad_f32 is 1.8x faster. It passed on one input, and a host timer printed a smaller number. On another run, the advantage shrinks. On a length that isn't divisible by 256, one output near the end changes.

That isn't a performance result yet. A trustworthy claim must show that both kernels computed the same operation, timed completed GPU work under identical conditions, and improved the scope users care about. Profilers explain where time went. They don't repair a broken comparison.

The Accelerator Architecture Field Guide established streaming multiprocessors, high-bandwidth memory, caches, and hardware ceilings. CUDA for ML Training established grids, blocks, asynchronous launches, and event timing. Now we turn those mechanics into an uncompromising measurement discipline built around five verification gates:

1. Gate 1: Reference Verification. Prove numerical output matches an independent double-precision CPU reference within declared tolerances ($|y_i - r_i| \leq \text{atol} + \text{rtol}|r_i|$), testing both power-of-two and arbitrary unaligned lengths.
2. Gate 2: Compute Sanitizer. Verify runtime memory safety and synchronization integrity (memcheck, racecheck, initcheck, synccheck), ensuring illegal memory accesses don't hide behind lucky allocations.
3. Gate 3: Clean Synchronized Timing. Measure steady-state GPU elapsed time using stream-bound CUDA events, warmups, and explicit L2 cache invalidation, isolating clean runs from profiler overhead.
4. Gate 4: Hardware Roofline Bound. Compare measured memory traffic and compute throughput against the theoretical hardware ceilings, verifying the kernel doesn't violate physical bandwidth limits.
5. Gate 5: Reproducible Receipt. Capture hardware state, GPU UUID, clocks, compiler flags, and raw sample hashes in an immutable record so any peer can reproduce the result.

One small kernel will carry the investigation. For every index $i$, it computes a floating-point triad:

$$y_i = \alpha x_i + z_i$$

It looks too small to fail. That makes it useful: if the evidence is weak here, a fused attention or quantization kernel will hide the same mistakes behind far more code.

Turn “faster” into a testable claim

Suppose a candidate replaces an existing triad_f32 implementation. Before opening a profiler, write the claim so another engineer could disprove it. The starter supplies a baseline and a deliberately incorrect variant, not two valid implementations with a measured speedup.

The odd length is intentional. With 256 threads per block, the last block contains threads whose indexes exceed $n-1$. A boundary guard must reject them. Testing only lengths divisible by block size leaves an off-by-one bug invisible.

The arithmetic also supplies an early hypothesis. One triad element performs one multiply and one add, counted as two floating-point operations (FLOPs). Its algorithmic minimum traffic is two float32 loads and one float32 store, or 12 bytes:

$$I_{algorithmic} = \frac{2\ \text{FLOPs}}{12\ \text{bytes}} \approx 0.167\ \text{FLOPs/byte}$$

That low arithmetic intensity suggests a bandwidth-sensitive kernel. It doesn't prove measured device-memory traffic equals 12 bytes per element, nor that high-bandwidth memory (HBM) is the current limit. Caches, transaction efficiency, launch overhead, and extra instructions can all move the measured point.

💡 Key insight: A performance hypothesis comes from operation and byte counts. A performance claim comes from a correctness-checked measurement on named hardware.

A cache-hot run processes 1,000,003 elements in 8.0 microseconds. Using the 12-byte algorithmic floor, what useful bandwidth does that imply, and why doesn't it prove HBM delivered that rate?

Answer

The calculation is $\frac{12 \times 1{,}000{,}003}{8.0 \times 10^{-6}}$, or about $1{,}500$ GB/s. That number counts algorithmic input and output bytes. It isn't a device-memory traffic counter, so cache hits can make useful bandwidth exceed the GPU's HBM specification without violating a hardware limit.

Gate 1: Reference verification and numerical tolerance

A trusted reference should be simpler than the candidate and independent enough to catch its mistakes. For triad, compute on the CPU in double, cast to float, and compare every element. A fused multiply-add (FMA) rounds the multiplication and addition together rather than rounding their intermediate product. The CPU reference rounds in double and then converts to float; the GPU's fmaf rounds directly to float. Higher precision helps here, but two-stage rounding isn't a universal bitwise oracle.^{[1]Reference 1Floating Point and IEEE 754 Compliance for NVIDIA GPUshttps://docs.nvidia.com/cuda/floating-point/index.html}

A checksum is too weak: two wrong values can cancel, and a single out-of-bounds write may miss sampled positions. Initialize output to a recognizable poison value before checking, so a skipped store can't inherit a previous correct result. Allocate each test's logical length, not a larger buffer that hides its tail overrun.

Floating-point equality needs a numerical contract. PyTorch's allclose rule is a useful statement of it:^{[2]Reference 2torch.allclosehttps://docs.pytorch.org/docs/stable/generated/torch.allclose}

$$|y_i - r_i| \leq \text{atol} + \text{rtol}|r_i|$$

Here $y_i$ is the candidate output and $r_i$ is the reference. Relative tolerance scales with nonzero reference magnitude. Absolute tolerance protects values near zero, where relative error alone can explode. Neither constant is universal. Choose them from dtype, operation count, expected rounding, input range, and downstream quality requirements, then freeze them before comparing candidates.

For bounded float32 inputs and one fused multiply-add, this lab uses atol = 2e-6 and rtol = 2e-5. Those are acceptance settings for this workload, not recommended defaults for arbitrary kernels. A reduction over thousands of values needs different analysis because accumulation order changes rounding.

Exercise the contract across shapes and values that expose distinct failure modes:

This CPU-only indexing exercise isolates the allocation trap. Python bounds checks stand in for the illegal access that a CUDA sanitizer should report; they don't simulate GPU execution. Predict why the same wrong guard behaves differently with spare allocation space.

Gate 2: Compute Sanitizer and runtime memory safety

Reference comparison catches wrong values that reach output. Runtime tools catch illegal execution that happens to leave checked values unchanged. NVIDIA Compute Sanitizer provides four specialized checkers:^{[3]Reference 3Compute Sanitizerhttps://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/}

• memcheck: Detects out-of-bounds global, shared, and local memory accesses, misaligned memory instructions, and illegal device pointers.
• racecheck: Detects shared-memory data hazards (RAW, WAR, WAW hazards where threads in a block access the same shared memory location without synchronization).
• initcheck: Flags reads of uninitialized global memory, catching kernels that assume memory allocations come pre-zeroed.
• synccheck: Catches illegal barrier synchronization, such as calling __syncthreads() inside divergent branches where not all warps arrive.

Run memcheck first. NVIDIA specifically recommends running memcheck before tools such as racecheck and synccheck, because race checking doesn't validate underlying pointer validity.^{[3]Reference 3Compute Sanitizerhttps://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/} Always configure --padding 32 to detect memory overruns that would otherwise land safely inside allocator alignment padding, and pass --error-exitcode 99 so automated test runners fail immediately when errors occur.

This sequence prevents an embarrassing trap. A profiler can report record-breaking memory throughput for a kernel that reads past buffer bounds. High hardware utilization never upgrades undefined behavior into a valid optimization.

Profiler scope: Nsight Systems versus Nsight Compute

Start wide enough to see ownership of delay, then zoom into one proven hotspot. NVIDIA separates those scopes across Nsight Systems and Nsight Compute.^{[4]Reference 4Nsight Systems User Guidehttps://docs.nvidia.com/nsight-systems/UserGuide/[5]Reference 5Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

Nsight Systems finds system-level bottlenecks

Open a system trace before assuming the hottest-looking kernel owns application delay. Nsight Systems attaches lightweight runtime hooks (CUPTI and OS scheduler tracing) with low overhead (typically 1% to 5%). Look for large host-to-device transfers, CPU preparation gaps, stream serialization, blocking scalar reads, and GPU idle bubbles. A long kernel is actionable only if it lies on the critical path for the claim.

For this lab, event synchronization after every sample intentionally serializes launches. Nsight Systems reveals that pattern immediately. It's correct for the isolated kernel-latency question, but it would distort a throughput experiment that expects multiple kernels or requests in flight. Scope determines whether synchronization is a measurement boundary or an application bottleneck.

Nsight Compute diagnoses kernel microarchitecture

Once triad_f32 is a known hotspot on the critical path, collect kernel metrics. Nsight Compute hooks directly into hardware performance monitoring units (PMUs) on the streaming multiprocessors. Its report sections cover launch configuration, memory workload analysis, GPU Speed of Light, and Roofline analysis.^{[5]Reference 5Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

Its baseline feature compares two kernel runs side-by-side. However, microarchitectural profiling changes the physical execution conditions.

Profiler distortions: multi-pass replay and counter perturbation

Hardware performance counters on GPU streaming multiprocessors are physically limited. An SM can't record hundreds of architectural counters in a single execution pass. To assemble a comprehensive report across instructions, cache hierarchies, and pipeline stalls, Nsight Compute uses multi-pass kernel replay.

During replay, the profiler saves the GPU memory state and re-executes the exact same kernel dozens of times under different hardware counter configurations. By default, Nsight Compute also enforces --cache-control all, explicitly flushing L1 and L2 caches before every pass to ensure deterministic counter collection.

This mechanism introduces massive execution distortion:

• A 15-microsecond kernel can take hundreds of milliseconds to execute under profiling.
• Cache lines that would remain resident in a tight loop get evicted before each replay pass.
• Warp scheduling interleaving and memory controller queues get serialized to capture fine-grained stalls.

Never quote elapsed time from an instrumented profiler report as a clean benchmark. Profilers explain why code behaves as it does; clean, un-instrumented benchmark runs measure how fast it actually finishes.

Sanitizers run outside timing

Compute Sanitizer instruments memory instructions and barrier synchronizations, adding substantial overhead and serializing concurrent warps. A clean sanitizer report is correctness evidence, not a latency sample. Run sanitizers to verify safety, then close them completely before taking benchmark measurements.

⚠️ Common mistake: Running every tool simultaneously produces an expensive, distorted execution whose timing answers no valid question. Keep correctness verification, microarchitectural diagnosis, and clean benchmark timing as separate, linked runs.

Gate 3: Synchronized timing and benchmark state control

“Kernel time” still leaves choices. State whether the product reuses the same resident data, streams new buffers, overlaps launches, or includes transfers. Different scopes can all be legitimate, but they can't share one unlabeled number.

Host versus device asynchrony: timing streams, not launch queues

CUDA kernel launches (<<<grid, block, shared, stream>>> or cudaLaunchKernel) are non-blocking asynchronous calls. When host code launches a kernel, the CPU constructs a launch descriptor packet, writes it into the driver's ring buffer (push-buffer in pinned memory), and returns control to the CPU thread immediately. This handoff typically takes 2 to 5 microseconds.

If an engineer wraps a host timer (std::chrono::high_resolution_clock or Python's time.perf_counter()) around a kernel launch without stream synchronization, the timer stops the instant the CPU finishes enqueuing the packet. The GPU hasn't even scheduled the blocks on SMs! The resulting "speedup" is merely a measurement of how fast the CPU pushed a descriptor into a queue.

CUDA events solve this by placing timestamp markers into the execution stream:

Here is what happens in hardware:

1. cudaEventRecord(start, stream) writes an event marker into the stream FIFO.
2. When the GPU execution engine reaches that marker, it records the hardware clock timestamp into the event object.
3. The kernel executes on the SMs.
4. When the kernel completes all thread blocks, the GPU reaches cudaEventRecord(stop, stream) and records the stop timestamp.
5. cudaEventSynchronize(stop) halts the host thread until the GPU signals that the stop event has retired.
6. cudaEventElapsedTime(&elapsed_ms, start, stop) calculates the true elapsed time directly from the hardware timestamps.

Coarse device synchronization (cudaDeviceSynchronize()) stalls every stream on the GPU, draining concurrent queues and killing pipeline overlap. Stream-specific event synchronization cleanly measures the targeted boundary without penalizing independent streams.

For tiny sub-microsecond kernels, time a batch of launches and divide by iteration count to amortize event overhead, keeping the batch size identical across comparisons. CUDA events provide roughly 0.5-microsecond resolution; zero or sub-microsecond measurements signal an insufficient measurement window rather than infinite speed.^{[6]Reference 6CUDA Runtime API: Event Managementhttps://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EVENT.html}

L2 cache invalidation: preventing residency bias

Modern GPU cache hierarchies feature substantial L2 caches:

• NVIDIA A100: 40 MB L2 cache
• NVIDIA H100: 50 MB L2 cache
• NVIDIA RTX 4090: 72 MB L2 cache
• NVIDIA B200: 60 MB L2 cache

Now consider the triad working set for $n = 1{,}000{,}003$ single-precision floats. Each buffer ($X$, $Z$, $Y$) consumes $4$ MB. Three buffers require $12$ MB total.

That 12 MB working set fits entirely inside the GPU's L2 cache!

In a naive benchmark loop that repeatedly launches on the same buffer addresses:

• Launch 1 fetches data from HBM3 (yielding ~2.0 to 3.35 TB/s).
• Launches 2 through 200 find every single cache line already resident in L2!
• L2 cache delivers up to 12 TB/s of aggregate bandwidth on an H100.
• The benchmark prints an elapsed time of ~3 microseconds, implying an impossible $4{,}000$ GB/s of bandwidth.

An engineer rejoices: "Our kernel hit 4 TB/s on H100!" In reality, the kernel never touched high-bandwidth memory. It merely benchmarked L2 cache latency.

In production LLM serving or training pipelines, incoming user prompts and fresh activations stream from DRAM. They don't arrive pre-warmed in L2 cache. Benchmarking an un-rotated small buffer creates an artificial residency bias that vanishes in production.

Two techniques prevent L2 residency bias:

1. Buffer Rotation: Allocate an array of $S$ distinct buffer triplets such that total memory footprint exceeds $2\times$ the L2 cache capacity (e.g. $S \times 12\ \text{MB} \geq 100\ \text{MB}$). Cycle through slot = i % slots on every iteration, forcing every launch to fetch from DRAM.
2. Cache Invalidation Sweep: If memory constraints prevent allocating multiple buffers, launch an explicit cache-clearing sweep kernel over a separate dummy buffer sized to $2\times$ L2 capacity between timed runs, evicting all cached lines back to DRAM.

Warm until mechanism is steady

Early iterations absorb CUDA context initialization, memory allocator pool growth, library loading, just-in-time compilation, page table mapping, and GPU clock frequency ramping.

Warmup iterations aren't arbitrary. Monitor latency, clocks, and temperature until they stabilize, then freeze the warmup count. Never discard slow samples that occur later in the run: thermal or power throttling during sustained execution is a physical hardware constraint, not an outlier.

Record clocks, power, and thermal state

GPU clocks dynamically adjust based on thermal headroom, power limits, and concurrent workloads. A candidate kernel measured after the card has heated up can appear slower even when its instruction stream is superior. NVIDIA System Management Interface (nvidia-smi) provides scriptable state queries:^{[7]Reference 7nvidia-smi documentationhttps://docs.nvidia.com/deploy/nvidia-smi/index.html}

On dedicated hardware with administrator privileges, lock clocks during benchmarking to eliminate frequency jitter:

On shared cloud VMs where clock locking is restricted, query and log the actual observed clocks before and after every benchmark round. Ensure both baseline and candidate run within the same clock and thermal envelope.

Preserve the distribution with alternating paired rounds (ABBA)

Reporting only the minimum execution time rewards random noise. Reporting a single mean obscures bimodal behavior and tail latency. Preserve raw samples, reporting sample count, median, p10, p90, and the p10-to-p90 spread.

To prevent thermal drift or background host activity from favoring one implementation, alternate rounds in an A, B, B, A sequence. Within each round $r$, compute the paired speedup:

$$\text{speedup}_r = \frac{\text{baseline median latency in round }r}{\text{candidate median latency in round }r}$$

Then evaluate the distribution of round-level speedups. If a candidate wins round 1 ($1.2\times$) but loses round 2 ($0.9\times$), pooling all samples into a single average hides that the candidate regressed once the GPU reached thermal equilibrium.

Two controlled rounds produce baseline and candidate medians of (12, 10) and (9, 10) microseconds. What are the round speedups, and why is one pooled speedup inadequate?

Answer

The round speedups are $12/10 = 1.2\times$ and $9/10 = 0.9\times$. The candidate wins one round and loses the other, so pooling samples into one number would hide a reversal that may come from drift or an unstable effect. Report both round-level results, inspect clocks and temperature, and collect more alternating rounds before making a speed claim.

Gate 4: Hardware Roofline and Speed-of-Light bounds

The Roofline model bounds attainable performance by the smaller of the compute ceiling and the bandwidth ceiling scaled by arithmetic intensity:^{[8]Reference 8Roofline: An Insightful Visual Performance Model for Multicore Architectureshttps://doi.org/10.1145/1498765.1498785}

$$P \leq \min(P_{peak},\ I \times BW)$$

Nsight Compute plots the achieved point against precision- and device-specific ceilings. The sloped region represents the memory-bandwidth bound; the flat region represents the peak-compute bound. The ridge is their intersection.^{[5]Reference 5Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

For triad, algorithmic intensity begins near 0.167 FLOPs/byte. Compare byte counts and bandwidth at the same memory level: an L2 hit is still an input load, but not an HBM read. More than 12 HBM bytes per element can indicate transaction inefficiency or redundant traffic; fewer can indicate cache reuse or another traffic-reducing mechanism. Under a genuinely streaming regime, approaching both the useful-byte floor and the bandwidth roof leaves little for shared-memory staging to remove.

Read several sections together:

One percentage isn't diagnosis. Low compute utilization can mean memory limit, too little parallel work, launch overhead, dependencies, or host starvation. High device-memory utilization can coexist with wasteful transactions. Roofline locates achieved point; system trace and detailed counters explain why it landed there.

For streaming triad, test this hypothesis:

1. Reference and sanitizer pass.
2. Nsight Systems confirms kernel dominates selected scope rather than copies or host gaps.
3. Nsight Compute shows low arithmetic intensity and strong device-memory pressure.
4. Measured bytes per element are compared with 12-byte algorithmic floor.
5. Optimization targets excess traffic or access order. Shared memory is rejected when it can't remove compulsory traffic.

If the measurements don't support that sequence, revise the hypothesis before changing the kernel.

Nsight Compute places triad_f32 near the sloped bandwidth roof, but Nsight Systems shows that the GPU is idle for most of the request while the CPU prepares inputs. What can you conclude, and which scope should you optimize first?

Answer

You can conclude that the isolated kernel is probably bandwidth-limited under the profiled workload. You can't conclude that kernel tuning will materially improve request latency. The system trace identifies CPU preparation on the request's critical path, so measure and repair that wider bottleneck before spending effort on a kernel already near its roof.

Run the triad investigation: five gates in code

The CUDA lab requires an NVIDIA GPU, a compatible driver/toolkit, and the profiling tools named below. It wasn't compiled or run on CUDA during this review; the displayed output is an acceptance shape, not a hardware result. CPU-only checks can validate the arithmetic, host helper functions, and receipt parser, but can't establish device correctness or speed.

The file includes buggy and fixed kernels, deterministic inputs, a CPU reference, event timing, same-buffer and buffer-rotated modes, and raw-sample CSV output. Read it in three parts: kernels and checking helpers, setup and test modes, then the timed loop. --profile launches one fixed kernel as an isolated diagnostic, not a warmed benchmark.

Compile with optimization and line information. Line information lets sanitizer and profiler reports point back to source without using debug build as performance binary:

A correct run prints the target device and a complete elementwise check. The output below describes the required shape, not a recorded CUDA run. Inspect the actual error values rather than requiring the zeros shown here. This fixed-size starter isn't a general allocator or launch planner; larger shapes need device-limit and integer-overflow checks too.

Now run the deliberately wrong i <= n guard under memcheck. --error-exitcode 99 makes a detected sanitizer error fail automation even when the target process exits cleanly, and pipefail keeps tee from hiding that status:

Because the grid rounds up and $n=1{,}000{,}003$, a thread with i == n executes. The report should identify an invalid global access in triad_buggy. Padding improves detection when adjacent allocations might otherwise hide an overrun. Without the sanitizer, the launch may appear to complete. Use the fixed mode with i < n, then require a clean report.^{[3]Reference 3Compute Sanitizerhttps://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/}

Run initcheck when a candidate changes which global data is initialized or read; it doesn't require shared memory. Add racecheck for shared-memory communication and synccheck for synchronization changes. None replaces memcheck or numerical comparison.

Capture the system timeline next. Timings printed while the profiler is attached are diagnostic only:

The full trace includes allocations, setup copies, the initial correctness readback, warmup, 200 timed launches, and final per-slot verification copies. Only the event-bounded sample region claims to exclude copies. Inside it, expect one launch and one event wait per sample; investigate long host gaps rather than mistaking setup or verification for timed work.

List the installed sections before collecting a kernel profile. Section sets vary by installation; don't assume a set named roofline exists. The command names four standard sections and makes cache flushing explicit. --clock-control none leaves existing clock policy unchanged.^{[9]Reference 9Nsight Compute CLIhttps://docs.nvidia.com/nsight-compute/NsightComputeCli/}

On a locked-down host, ncu may return ERR_NVGPUCTRPERM. Counter-access policy blocked metric collection; the kernel didn't fail. Use an approved profiling host or ask an administrator to enable performance counters rather than changing privilege policy on a shared machine.^{[5]Reference 5Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

Read the achieved point, device-memory bytes, cache behavior, and SpeedOfLight sections together. Compare measured traffic per element with the 12-byte algorithmic count. Useful bandwidth above the device's HBM specification can reflect data served by cache, rather than impossible hardware performance.

This isolated, cache-flushed profile doesn't establish the residency of the clean same-buffer benchmark. Nsight Compute normally flushes caches between replay passes and may control clocks; replay mode also changes execution. For a cache-primed application, a separately designed experiment can use application replay with --cache-control none so its setup is repeated for each pass. Record these choices and never mix a cold profile's byte count with a hot run's elapsed time in one Roofline point.^{[5]Reference 5Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

Finally collect clean timing outside tools:

The starter warms for at least 20 launches and enough launches to touch every rotation slot twice. Treat that as an initial policy, not proof of steady state. Increase it if latency, clocks, temperature, or cache counters haven't stabilized on the target machine.

The post-timing check validates each slot's final output, not every intermediate launch. Keep adversarial correctness runs separate so an intermittent omitted store can't hide behind a previous correct value. Record nvcc, nsys, ncu, and compute-sanitizer versions alongside the exact commands; profiler section availability and supported metrics vary by tool and GPU.

--bench-hot is a convenient flag name, not evidence of cache residency: the program reports it as same-buffer. Percentiles use linear interpolation at index $q(N-1)$, so an even-count median averages the middle pair. Samples remain in acquisition order in the CSV. Match the selected CUDA device to its UUID, especially when CUDA_VISIBLE_DEVICES remaps ordinals; a GPU name alone isn't enough on a multi-GPU host.

Don't compare a hot baseline with a rotated candidate. Run both versions in the same regime and alternate order across rounds.

Diagnose contradictory evidence

GPU investigations often return signals that look inconsistent. Resolve by asking which scope each signal describes.

Failure diagnosis should end with one controlled experiment. Changing block size, vector width, memory layout, and dtype together destroys attribution even if final number improves.

Gate 5: Reproducible receipts and machine-readable evidence

A chat message saying “A is faster” isn't reproducible. Save a machine-readable receipt plus raw artifacts. At minimum, another engineer should know exactly what ran, where it ran, what passed, and which files support the claim.

Store reports with the source revision, not in a personal profiler workspace that disappears. Hash the source and raw sample files. If the workload contains proprietary input, archive the deterministic generator, shape distribution, and a redacted data fingerprint so the receipt remains useful without leaking data.

The Python fixture below reads real temporary CSV files containing synthetic timings. It checks saved-byte hashes, sample order, positive finite durations, and matching comparison conditions before computing one round's ratio. It deliberately doesn't print “claim ready”: artifact integrity and a favorable ratio don't prove correctness, representative hardware conditions, or a repeatable improvement.

A hash binds a receipt to bytes; it doesn't prove those bytes came from a GPU. Collect hardware state, source identity, correctness logs, sanitizer exit status, and timing artifacts through a trusted harness. This small parser checks only timing-file integrity and declared conditions, not that entire collection process. Run it per matched round, then examine the round-level distribution.

The five verification gates in production review

Before submitting or merging any GPU optimization patch, run through the complete verification sequence:

1. Gate 1: Reference Verification. Test numerical equivalence against a higher-precision reference on power-of-two, prime, and unaligned lengths ($n = 0, 1, 255, 256, 257, 1{,}000{,}003$) with explicit poison patterns.
2. Gate 2: Compute Sanitizer. Run compute-sanitizer --tool memcheck --padding 32 --error-exitcode 99 to prove that no out-of-bounds reads or writes occur on any exercised path.
3. Gate 3: Synchronized Timing. Time completed GPU execution via stream-bound CUDA events (cudaEventRecord), warm the card past initial JIT and driver overhead, invalidate L2 cache lines across iterations, and run alternating ABBA rounds.
4. Gate 4: Hardware Roofline Bound. Calculate arithmetic intensity and verify that measured memory bandwidth doesn't violate physical HBM limits. If bandwidth looks impossibly high, check for unevicted L2 cache hits.
5. Gate 5: Reproducible Receipt. Persist the git commit SHA, binary hash, GPU UUID, driver version, clocks, power limits, raw sample CSVs, and sanitizer output in a structured receipt.json.

Optimization earns trust when the faster candidate remains the same program for approved inputs and the evidence survives another engineer's rerun. Kernel engineering starts after that foundation, not before it.

Mastery check

Evaluation rubric

• Foundational: Separate reference checking, sanitizer evidence, profiler diagnosis, and clean timing into distinct runs.
• Intermediate: Compute useful bandwidth from declared algorithmic bytes while distinguishing it from measured HBM traffic.
• Advanced: Design alternating benchmark rounds, interpret system and kernel profiles at their proper scopes, and defend a speedup from a reproducible evidence receipt.

Common pitfalls

• Treating a clean sampled output as proof that an illegal memory access is harmless.
• Comparing a cache-hot candidate with a buffer-rotated baseline.
• Publishing profiler duration as clean latency or pooling samples across drifting rounds.

Follow-up questions

Why can an i <= n kernel pass an elementwise check for n=257 when its allocation has 512 elements?

Answer

It writes all 257 logical outputs correctly and also writes element 257. A comparison restricted to the logical output misses that extra store, while a memory checker sees an allocation-valid access. Use exact-sized allocations, output poisoning, and sanitizer padding to expose different parts of the failure.

A warm same-buffer run is fast, but a cache-flushed Nsight Compute report shows high HBM traffic. Can you divide that report's bytes by the warm event duration?

Answer

No. The numerator and denominator describe different cache regimes. Use traffic and duration from a consistent profiling experiment for its Roofline point, and report the clean benchmark separately. Record replay, cache, and clock controls so the difference is interpretable.