LearnInference & Production ScaleCorrect GPU Profiling and Benchmarking

🚀HardInference Optimization

Correct GPU Profiling and Benchmarking

Build trustworthy GPU performance evidence: prove kernel correctness, choose system or kernel profiling scope, control benchmark state, and preserve a reproducible receipt before claiming a speedup.

27 min read

Learning path

Step 145 of 192 in the full curriculum

Accelerator Architecture Field Guide GPU Kernel Performance Engineering Lab

Personalize this lesson

Adapt explanations and teaching visuals to your background and preferred voice.

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 the same 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. Keep those mechanics. Now turn them into a measurement discipline.

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 the candidate triad_f32 replaces a scalar baseline. Before opening a profiler, write the claim so another engineer could disprove it.

Claim field	Concrete contract for the running kernel
Semantics	For each valid index, return `fmaf(alpha, x[i], z[i])` within declared tolerance
Workload	`float32`, $n=1{,}000{,}003$ , deterministic bounded inputs
Scope	One completed kernel launch, excluding allocation and host-to-device copies
Cache regime	Report cache-hot and buffer-rotated measurements separately
Baseline	Same compiler flags, device, inputs, output check, and timing method; record each version's launch shape
Environment	GPU UUID, driver, CUDA runtime, source revision, clocks, power limit, and temperature
Decision	Keep candidate only if correctness gates pass and measured scope improves

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 can leave 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_{min} = \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.

Prove output before studying speed

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 once to float, and compare every element. Both paths use a fused multiply-add (FMA), which rounds multiplication plus addition once rather than after each operation.^{[1]Reference 1CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/} A checksum is too weak: two wrong values can cancel, and a single out-of-bounds write may miss sampled positions.

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:

Test	Why it exists	Failure signature
$n=0$ and $n=1$	empty and smallest legal work	launch or guard assumption
$n=255,256,257$	both sides of one block boundary	last-block indexing bug
$n=1{,}000{,}003$	realistic rounded-up grid	rare tail corruption
zeros and cancellation	reference values near zero	missing absolute tolerance
largest lab magnitude, $10^{10}$	large-value rounding and cancellation	absolute- or relative-error spike
repeated seeds	varied values with reproducible inputs	input-sensitive mismatch

Reference comparison catches wrong values that reach output. Runtime tools catch illegal execution that happens to leave checked values unchanged. NVIDIA Compute Sanitizer is a functional correctness suite: memcheck detects out-of-bounds and misaligned memory accesses, racecheck detects shared-memory hazards, initcheck detects uninitialized device-memory reads, and synccheck detects invalid synchronization use.^{[3]Reference 3Compute Sanitizerhttps://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/}

Run memcheck first. NVIDIA specifically recommends that order before tools such as racecheck and synccheck, because those tools don't replace memory-access checking.^{[3]Reference 3Compute Sanitizerhttps://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/}

Diagram showing Candidate + fixed workload, Reference + tolerance gate, Sanitizers clean?, and No. — Candidate + fixed workload, Reference + tolerance gate, Sanitizers clean?, and No.

The order prevents a familiar trap. A profiler can report excellent memory throughput for a kernel that reads one element past the buffer. High utilization never upgrades undefined behavior into a valid optimization.

Choose instrument by question

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/}

Question	Instrument	Evidence it can supply	What it can't establish alone
Why is application waiting?	Nsight Systems	CPU scheduling, CUDA API calls, transfers, kernels, streams, synchronization, idle gaps	exact instruction or memory limit inside one kernel
Why is this kernel slow?	Nsight Compute	launch resources, SpeedOfLight sections, cache and device-memory traffic, warp stalls, Roofline point	user-visible end-to-end latency or output correctness
Is CUDA execution legal?	Compute Sanitizer	invalid access, shared-memory hazard, uninitialized read, invalid barrier use	representative performance under normal execution
How long did completed device work take?	CUDA events or a synchronizing benchmark harness	device elapsed time for declared event boundaries	cause of the elapsed time

Nsight Systems finds scope

Open a system trace before assuming the hottest-looking kernel owns the application delay. Look for large host-to-device copies, CPU input gaps, many tiny launches, stream serialization, blocking scalar reads, and idle GPU regions. 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 should reveal that pattern. It's correct for the isolated kernel-latency question, but it would distort a throughput experiment that expects several kernels or requests in flight. Scope determines whether synchronization is a measurement boundary or a workload bug.

Nsight Compute explains one kernel

Once triad_f32 is a known hotspot, collect kernel metrics. Nsight Compute's report sections cover launch configuration, memory workload, GPU Speed of Light, and Roofline analysis. Its baseline feature can compare reports, but metric collection and kernel replay add overhead.^{[5]Reference 5Nsight Compute Profiling Guidehttps://docs.nvidia.com/nsight-compute/ProfilingGuide/}

Never quote elapsed time from a heavily instrumented profiler run as a clean benchmark. Profile to explain. Rerun outside the profiler to measure.

Sanitizers run outside timing

Compute Sanitizer changes execution cost and may serialize behavior. A clean sanitizer report is correctness evidence, not latency sample. Preserve its command, tool version, exit status, and report alongside performance artifacts.

⚠️ Common mistake: Running every tool at once produces one expensive, perturbed execution whose timing answers no clean question. Keep correctness, diagnosis, and measurement as linked but separate runs.

Read Roofline and SpeedOfLight without guessing

The Roofline model bounds attainable performance by the smaller of the compute ceiling and the bandwidth ceiling scaled by arithmetic intensity:^{[6]Reference 6Roofline: 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. If the profiler reports substantially more than 12 bytes of device-memory traffic per element, the measured intensity is even lower than the algorithmic estimate. That gap points toward transaction inefficiency or redundant traffic. If measured traffic stays near the compulsory minimum and the point approaches the bandwidth roof, adding shared memory probably can't remove two input reads and one output write. It may add instructions instead.

Read several sections together:

Nsight Compute evidence	Question answered	Sound next move
`LaunchStats`	Did grid, block, registers, or shared memory constrain launch?	change launch resources only when metric identifies pressure
`SpeedOfLight`	Which compute or memory paths approach throughput ceiling?	inspect saturated path and corroborating stalls
`MemoryWorkloadAnalysis`	How much traffic reaches cache levels and device memory?	compare measured bytes with algorithmic minimum
Roofline chart	Does achieved point sit in bandwidth or compute region?	target bytes/reuse on slope; target math pipeline near flat roof

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, likely sequence is:

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

That conclusion may sound uneventful. Avoiding a useless rewrite saves engineering time and keeps risk out of the hot path.

Design benchmark around real regime

“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.

Measurement question	Include	Exclude	Timing method
isolated launch latency	one kernel on resident buffers	allocation and copies	CUDA events around launch
steady pipeline throughput	representative sequence with normal overlap	one-time setup	events around batch or end-to-end synchronized wall clock
first request latency	compile, lazy initialization, allocation, and transfers that user experiences	unrelated process startup	synchronized end-to-end clock
cache-sensitive stream	rotated working set sized beyond relevant cache	accidental same-buffer reuse	events plus cache counters

Synchronize boundaries, not every operation

CUDA launches are asynchronous with respect to the host. A CPU timer can stop after enqueue while the device still works. CUDA events record progress in a stream and can measure elapsed device time after the end event completes.^{[7]Reference 7CUDA semantics.https://docs.pytorch.org/docs/2.13/notes/cuda.html}

For one launch:

Finish unrelated setup before start event.
Record start event in same stream as kernel.
Launch kernel.
Record end event.
Wait for end event, then read elapsed time.

A device-wide cudaDeviceSynchronize() inside the timed region adds a boundary that the production path may not contain. Use an event dependency when it matches the question. For tiny kernels, time a batch of launches and divide when event overhead is large relative to the work, while preserving the same batch size for the baseline and candidate.

PyTorch's benchmark utilities address accelerator synchronization and warmup for similar reasons.^{[8]Reference 8PyTorch Benchmarkhttps://docs.pytorch.org/tutorials/recipes/recipes/benchmark.html} A harness is helpful, but its defaults still need to match intended scope.

Warm until mechanism is steady

Early iterations can include CUDA context creation, allocator growth, library initialization, just-in-time compilation, page mapping, cache fill, and clock ramp. Warmup count isn't magic. Watch latency, clocks, temperature, and selected algorithm until they stabilize, then record count and stopping rule.

Don't silently discard inconvenient later samples. Thermal or power throttling after warmup is part of the environment, not an outlier by definition.

Name cache policy

Reusing one buffer can make data cache-resident. Rotating across a working set larger than the relevant cache better represents streaming data. Neither regime is universally honest.

Schematic normalized latency over six observations for three GPU benchmark regimes. An un-warmed run starts high and settles, a same-buffer run stays lower when its working set remains cache-resident, and warmed rotating buffers remain near the intended steady-state level. The values compare measurement regimes, not hardware results. — First-use work inflates early samples, while one reused buffer can create a different cache-hot workload. Warmup removes setup bias; buffer rotation changes data-residency regime. Pick regime from product path and label it.

The plotted values are schematic, not benchmark results. In an actual run, the total rotated footprint should exceed the target cache by a deliberate margin, and Nsight Compute cache metrics should confirm the assumed regime. Allocation size alone can't prove every access missed the cache.

Record clocks, power, and thermal state

GPU clocks respond to idle state, temperature, power limit, and other work. A candidate measured after card heats up can look slower even when its code is better. NVIDIA System Management Interface (nvidia-smi) exposes scriptable queries, supported clock controls, and power-limit state.^{[9]Reference 9nvidia-smi documentationhttps://docs.nvidia.com/deploy/nvidia-smi/index.html}

Capture the state before and after each comparison:

terminal

nvidia-smi --query-gpu=timestamp,uuid,name,driver_version,pstate,clocks.sm,clocks.mem,temperature.gpu,power.draw,power.limit --format=csv

On isolated, administrator-controlled hardware, supported clock locking can reduce variance. It requires privileges and differs by GPU. Users on a shared fleet should record observed clocks and power instead of changing machine policy. Either way, the baseline and candidate need the same power limit, thermal envelope, and competing-work policy.

Preserve distribution

Reporting only the minimum rewards noise. A single mean hides tails. Keep raw samples, then report at least the sample count, median, p10, p90, and p10-to-p90 spread. Compare the baseline and candidate in alternating rounds such as A, B, B, A so gradual clock or thermal drift doesn't always favor the same version.

Within each controlled round $r$ , compute speedup from like-for-like medians:

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

Then summarize the distribution of round-level speedups. State the number of rounds and a spread or interval instead of treating 200 adjacent launches on one thermal trajectory as 200 independent comparisons. Also report absolute time. A 2x speedup from 2 microseconds to 1 microsecond may disappear in an application dominated by launch, copies, or queueing.

Run the triad investigation

The lab uses one complete CUDA file. It includes buggy and fixed kernels, deterministic seeded inputs, a CPU reference, a tolerance check, event timing, cache-hot and buffer-rotated modes, and ordered raw-sample CSV output. --profile launches one fixed kernel so Nsight Compute doesn't collect hundreds of samples.

triad_bench.cu

#include <cuda_runtime.h>

#include <algorithm>
#include <cstddef>
#include <cmath>
#include <cstdlib>
#include <fstream>
#include <iomanip>
#include <iostream>
#include <string>
#include <utility>
#include <vector>

static void cuda_check(cudaError_t status, const char* call) {
    if (status != cudaSuccess) {
        std::cerr << call << ": " << cudaGetErrorString(status) << "\n";
        std::exit(1);
    }
}

#define CUDA_CHECK(call) cuda_check((call), #call)

__global__ void triad_buggy(
    const float* x,
    const float* z,
    float* y,
    std::size_t n,
    float alpha
) {
    const std::size_t i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i <= n) {
        y[i] = fmaf(alpha, x[i], z[i]);
    }
}

__global__ void triad_f32(
    const float* x,
    const float* z,
    float* y,
    std::size_t n,
    float alpha
) {
    const std::size_t i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        y[i] = fmaf(alpha, x[i], z[i]);
    }
}

static void launch_triad(
    bool buggy,
    const float* x,
    const float* z,
    float* y,
    std::size_t n,
    float alpha,
    cudaStream_t stream = nullptr
) {
    if (n == 0) {
        return;
    }
    constexpr int threads = 256;
    const int blocks = static_cast<int>((n + threads - 1) / threads);
    if (buggy) {
        triad_buggy<<<blocks, threads, 0, stream>>>(x, z, y, n, alpha);
    } else {
        triad_f32<<<blocks, threads, 0, stream>>>(x, z, y, n, alpha);
    }
    CUDA_CHECK(cudaGetLastError());
}

struct Verification {
    std::size_t failures;
    double max_abs;
    double max_rel;
};

static Verification verify(
    const std::vector<float>& x,
    const std::vector<float>& z,
    float* device_y,
    std::size_t n,
    float alpha
) {
    if (n == 0) {
        return {0, 0.0, 0.0};
    }

    std::vector<float> actual(n);
    CUDA_CHECK(cudaMemcpy(
        actual.data(),
        device_y,
        actual.size() * sizeof(float),
        cudaMemcpyDeviceToHost
    ));

    constexpr double atol = 2e-6;
    constexpr double rtol = 2e-5;
    Verification result{0, 0.0, 0.0};

    for (std::size_t i = 0; i < n; ++i) {
        const double reference64 = std::fma(
            static_cast<double>(alpha),
            static_cast<double>(x[i]),
            static_cast<double>(z[i])
        );
        const float reference = static_cast<float>(reference64);
        const double abs_error = std::abs(
            static_cast<double>(actual[i]) - static_cast<double>(reference)
        );
        const double rel_error = abs_error / std::max(std::abs(static_cast<double>(reference)), 1e-30);
        const double allowed = atol + rtol * std::abs(static_cast<double>(reference));

        result.max_abs = std::max(result.max_abs, abs_error);
        result.max_rel = std::max(result.max_rel, rel_error);
        if (!std::isfinite(actual[i]) || abs_error > allowed) {
            ++result.failures;
        }
    }
    return result;
}

static void fill_inputs(
    std::vector<float>& x,
    std::vector<float>& z,
    unsigned int seed,
    float alpha
) {
    for (std::size_t i = 0; i < x.size(); ++i) {
        const int x_code = static_cast<int>((i * 29 + seed * 17) % 251) - 125;
        const int z_code = static_cast<int>((i * 31 + seed * 13) % 239) - 119;
        x[i] = static_cast<float>(x_code) / 63.0f;
        z[i] = static_cast<float>(z_code) / 71.0f;
        if (i % 4096 == 0) {
            x[i] = 0.0f;
            z[i] = 0.0f;
        } else if (i % 4096 == 1) {
            x[i] = 1.0f;
            z[i] = -alpha;
        } else if (i % 4096 == 2) {
            x[i] = 1e10f;
            z[i] = -1e10f;
        }
    }
}

static double percentile(const std::vector<float>& sorted, double q) {
    const auto index = static_cast<std::size_t>(q * (sorted.size() - 1));
    return sorted[index];
}

static std::string cuda_version(int encoded) {
    return std::to_string(encoded / 1000) + "." + std::to_string((encoded % 1000) / 10);
}

int main(int argc, char** argv) {
    const std::string mode = argc > 1 ? argv[1] : "--check";
    const bool bench_hot = mode == "--bench-hot";
    const bool bench_rotate = mode == "--bench-rotate";
    const bool benchmark = bench_hot || bench_rotate;

    if (
        argc > 3 ||
        (
            mode != "--check" &&
            mode != "--buggy" &&
            mode != "--profile" &&
            !benchmark
        )
    ) {
        std::cerr << "usage: ./triad_bench [--check|--buggy|--profile|--bench-hot|--bench-rotate] [samples.csv]\n";
        return 2;
    }
    if (argc == 3 && !benchmark) {
        std::cerr << "a sample CSV path is valid only for benchmark modes\n";
        return 2;
    }

    constexpr std::size_t n = 1'000'003;
    constexpr float alpha = 1.25f;
    const std::size_t bytes = n * sizeof(float);

    cudaDeviceProp properties{};
    CUDA_CHECK(cudaGetDeviceProperties(&properties, 0));
    int runtime_version = 0;
    int driver_api_version = 0;
    CUDA_CHECK(cudaRuntimeGetVersion(&runtime_version));
    CUDA_CHECK(cudaDriverGetVersion(&driver_api_version));
    std::cout << "device=" << properties.name
              << " compute_capability=" << properties.major << "." << properties.minor
              << " l2_bytes=" << properties.l2CacheSize
              << " cuda_runtime=" << cuda_version(runtime_version)
              << " cuda_driver_api=" << cuda_version(driver_api_version) << "\n";

    std::vector<float> host_x(n);
    std::vector<float> host_z(n);
    fill_inputs(host_x, host_z, 7, alpha);

    int slots = 1;
    if (bench_rotate) {
        const std::size_t slot_bytes = 3 * bytes;
        const std::size_t target_bytes = std::max(
            8 * slot_bytes,
            2 * static_cast<std::size_t>(properties.l2CacheSize) + slot_bytes
        );
        const std::size_t desired_slots = (target_bytes + slot_bytes - 1) / slot_bytes;

        std::size_t free_bytes = 0;
        std::size_t total_bytes = 0;
        CUDA_CHECK(cudaMemGetInfo(&free_bytes, &total_bytes));
        const std::size_t affordable_slots = std::max<std::size_t>(1, free_bytes / (4 * slot_bytes));
        if (desired_slots > affordable_slots) {
            std::cerr << "buffer rotation needs " << desired_slots * slot_bytes
                      << " bytes but conservative allocation budget is "
                      << affordable_slots * slot_bytes << " bytes\n";
            return 1;
        }
        slots = static_cast<int>(desired_slots);
        std::cout << "rotation_target_bytes=" << target_bytes
                  << " rotation_working_set_bytes=" << slots * slot_bytes
                  << " free_bytes=" << free_bytes
                  << " total_bytes=" << total_bytes << "\n";
    }

    std::vector<float*> xs(slots, nullptr);
    std::vector<float*> zs(slots, nullptr);
    std::vector<float*> ys(slots, nullptr);

    for (int slot = 0; slot < slots; ++slot) {
        CUDA_CHECK(cudaMalloc(reinterpret_cast<void**>(&xs[slot]), bytes));
        CUDA_CHECK(cudaMalloc(reinterpret_cast<void**>(&zs[slot]), bytes));
        CUDA_CHECK(cudaMalloc(reinterpret_cast<void**>(&ys[slot]), bytes));
        CUDA_CHECK(cudaMemcpy(xs[slot], host_x.data(), bytes, cudaMemcpyHostToDevice));
        CUDA_CHECK(cudaMemcpy(zs[slot], host_z.data(), bytes, cudaMemcpyHostToDevice));
    }

    if (mode == "--buggy") {
        launch_triad(true, xs[0], zs[0], ys[0], n, alpha);
        CUDA_CHECK(cudaDeviceSynchronize());
        std::cout << "buggy launch completed; inspect it under Compute Sanitizer\n";
    } else if (mode == "--profile") {
        launch_triad(false, xs[0], zs[0], ys[0], n, alpha);
        CUDA_CHECK(cudaDeviceSynchronize());
        std::cout << "profile launch complete n=" << n << "\n";
    } else if (mode == "--check") {
        const std::vector<std::pair<std::size_t, unsigned int>> cases = {
            {0, 7}, {1, 7}, {255, 7}, {256, 7}, {257, 7},
            {n, 7}, {n, 19}, {n, 101},
        };
        for (const auto& [test_n, seed] : cases) {
            fill_inputs(host_x, host_z, seed, alpha);
            CUDA_CHECK(cudaMemcpy(xs[0], host_x.data(), bytes, cudaMemcpyHostToDevice));
            CUDA_CHECK(cudaMemcpy(zs[0], host_z.data(), bytes, cudaMemcpyHostToDevice));
            launch_triad(false, xs[0], zs[0], ys[0], test_n, alpha);
            CUDA_CHECK(cudaDeviceSynchronize());
            const Verification correctness = verify(host_x, host_z, ys[0], test_n, alpha);
            std::cout << std::scientific
                      << "seed=" << seed
                      << " n=" << test_n
                      << " correctness=" << (correctness.failures == 0 ? "PASS" : "FAIL")
                      << " failures=" << correctness.failures
                      << " max_abs=" << correctness.max_abs
                      << " max_rel=" << correctness.max_rel << "\n";
            if (correctness.failures != 0) {
                return 1;
            }
        }
    } else {
        launch_triad(false, xs[0], zs[0], ys[0], n, alpha);
        CUDA_CHECK(cudaDeviceSynchronize());
        const Verification correctness = verify(host_x, host_z, ys[0], n, alpha);
        std::cout << std::scientific
                  << "correctness=" << (correctness.failures == 0 ? "PASS" : "FAIL")
                  << " failures=" << correctness.failures
                  << " max_abs=" << correctness.max_abs
                  << " max_rel=" << correctness.max_rel << "\n";

        if (correctness.failures != 0) {
            return 1;
        }

        const int warmup = std::max(20, 2 * slots);
        constexpr int samples = 200;
        for (int i = 0; i < warmup; ++i) {
            const int slot = i % slots;
            launch_triad(false, xs[slot], zs[slot], ys[slot], n, alpha);
        }
        CUDA_CHECK(cudaDeviceSynchronize());

        cudaEvent_t start{};
        cudaEvent_t stop{};
        CUDA_CHECK(cudaEventCreate(&start));
        CUDA_CHECK(cudaEventCreate(&stop));
        std::vector<float> sample_ms;
        sample_ms.reserve(samples);

        for (int i = 0; i < samples; ++i) {
            const int slot = i % slots;
            CUDA_CHECK(cudaEventRecord(start));
            launch_triad(false, xs[slot], zs[slot], ys[slot], n, alpha);
            CUDA_CHECK(cudaEventRecord(stop));
            CUDA_CHECK(cudaEventSynchronize(stop));
            float elapsed_ms = 0.0f;
            CUDA_CHECK(cudaEventElapsedTime(&elapsed_ms, start, stop));
            sample_ms.push_back(elapsed_ms);
        }

        const std::string sample_path = argc == 3
            ? argv[2]
            : (bench_hot ? "samples-hot.csv" : "samples-rotate.csv");
        std::ofstream sample_file(sample_path);
        if (!sample_file) {
            std::cerr << "could not open sample CSV: " << sample_path << "\n";
            return 1;
        }
        sample_file << "sample,slot,elapsed_us\n" << std::fixed << std::setprecision(6);
        for (int i = 0; i < samples; ++i) {
            sample_file << i << "," << i % slots << "," << sample_ms[i] * 1000.0 << "\n";
        }

        std::vector<float> sorted_ms = sample_ms;
        std::sort(sorted_ms.begin(), sorted_ms.end());
        const double p10_ms = percentile(sorted_ms, 0.10);
        const double median_ms = percentile(sorted_ms, 0.50);
        const double p90_ms = percentile(sorted_ms, 0.90);
        const double useful_gb_s = (3.0 * bytes) / (median_ms / 1000.0) / 1e9;

        std::cout << std::fixed << std::setprecision(3)
                  << "regime=" << (bench_hot ? "cache-hot" : "buffer-rotated")
                  << " slots=" << slots
                  << " warmup=" << warmup
                  << " samples=" << samples
                  << " raw_samples=" << sample_path << "\n"
                  << "p10_us=" << p10_ms * 1000.0
                  << " median_us=" << median_ms * 1000.0
                  << " p90_us=" << p90_ms * 1000.0 << "\n"
                  << "useful_gb_s=" << useful_gb_s
                  << " (12 algorithmic bytes per element)\n";

        CUDA_CHECK(cudaEventDestroy(start));
        CUDA_CHECK(cudaEventDestroy(stop));
    }

    for (int slot = 0; slot < slots; ++slot) {
        CUDA_CHECK(cudaFree(xs[slot]));
        CUDA_CHECK(cudaFree(zs[slot]));
        CUDA_CHECK(cudaFree(ys[slot]));
    }
    return 0;
}

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

terminal

mkdir -p receipts
nvcc -O3 -lineinfo -std=c++17 triad_bench.cu -o triad_bench
./triad_bench --check | tee receipts/correctness.txt

A correct run prints the target device and a complete elementwise check. The exact device name varies. These bounded inputs should normally agree bit for bit with the chosen fused reference, while the tolerance gate still records the acceptance contract:

Expected output shape

device=<your GPU> compute_capability=<major.minor> l2_bytes=<bytes> cuda_runtime=<major.minor> cuda_driver_api=<major.minor>
seed=7 n=0 correctness=PASS failures=0 max_abs=0.000000e+00 max_rel=0.000000e+00
seed=7 n=1 correctness=PASS failures=0 max_abs=0.000000e+00 max_rel=0.000000e+00
seed=7 n=255 correctness=PASS failures=0 max_abs=0.000000e+00 max_rel=0.000000e+00
seed=7 n=256 correctness=PASS failures=0 max_abs=0.000000e+00 max_rel=0.000000e+00
seed=7 n=257 correctness=PASS failures=0 max_abs=0.000000e+00 max_rel=0.000000e+00
seed=7 n=1000003 correctness=PASS failures=0 max_abs=0.000000e+00 max_rel=0.000000e+00
seed=19 n=1000003 correctness=PASS failures=0 max_abs=0.000000e+00 max_rel=0.000000e+00
seed=101 n=1000003 correctness=PASS failures=0 max_abs=0.000000e+00 max_rel=0.000000e+00

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:

terminal

set -o pipefail
compute-sanitizer --error-exitcode 99 --tool memcheck ./triad_bench --buggy 2>&1 \
  | tee receipts/memcheck-buggy.txt

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. Without the sanitizer, the launch may appear to complete. Replace the guard with i < n by using fixed mode, then require a clean report:

terminal

set -o pipefail
compute-sanitizer --error-exitcode 99 --tool memcheck ./triad_bench --check 2>&1 \
  | tee receipts/memcheck-fixed.txt

Required sanitizer ending

ERROR SUMMARY: 0 errors

If next candidate adds shared memory or barriers, extend gate with racecheck, initcheck, and synccheck. Current triad has neither shared-memory communication nor block barriers, so those tools shouldn't be added as ceremony.

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

terminal

nsys profile \
  --trace=cuda,osrt \
  --sample=none \
  --output=receipts/triad_systems \
  ./triad_bench --bench-rotate receipts/nsys-diagnostic-samples.csv

The trace should show warmup launches followed by one launch and one event wait per sample. Unexpected copies or long host gaps would contradict the isolated resident-buffer claim.

Collect one kernel profile with Roofline set:

terminal

ncu \
  --set roofline \
  --launch-count 1 \
  --export receipts/triad_compute \
  ./triad_bench --profile

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 achieved point, device-memory bytes, cache behavior, and SpeedOfLight sections together. Compare measured traffic per element with 12-byte algorithmic floor. A cache-hot run may report useful bandwidth above device HBM specification because useful-byte calculation counts algorithmic bytes while data is served by cache. That's evidence of different regime, not impossible hardware.

Finally collect clean timing outside tools:

terminal

nvidia-smi --query-gpu=timestamp,uuid,name,driver_version,pstate,clocks.sm,clocks.mem,temperature.gpu,power.draw,power.limit --format=csv \
  | tee receipts/gpu-state-before.csv

./triad_bench --bench-hot receipts/bench-hot-samples.csv | tee receipts/bench-hot.txt
./triad_bench --bench-rotate receipts/bench-rotate-samples.csv | tee receipts/bench-rotate.txt

nvidia-smi --query-gpu=timestamp,uuid,name,driver_version,pstate,clocks.sm,clocks.mem,temperature.gpu,power.draw,power.limit --format=csv \
  | tee receipts/gpu-state-after.csv

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.

Don't compare hot baseline with rotated candidate. Run both versions in 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.

Observation	Tempting conclusion	Better diagnosis	Next check
host timer is tiny	kernel is fast	timer stopped after enqueue	CUDA events or synchronized boundary
sanitizer fails but sampled outputs match	error is harmless	illegal access missed sampled state	reject candidate and repair bounds
`nvidia-smi` utilization is high	kernel is efficient	GPU stayed busy doing some work	Nsight Systems critical path, then kernel counters
SpeedOfLight memory percentage is high	no optimization remains	one memory path is busy, perhaps wastefully	bytes per element and transaction efficiency
cache-hot timing wins	candidate is universally faster	benchmark changed residency regime	same-buffer versus rotated A/B under same policy
profiler duration is slower	optimization regressed	instrumentation or replay perturbed execution	clean rerun outside profiler
median improves while p90 worsens	clear win	distribution changed or throttling emerged	raw samples, clocks, power, temperature, paired order
kernel improves but request latency doesn't	profiler is wrong	optimized scope isn't end-to-end bottleneck	return to Nsight Systems and critical path

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.

Preserve evidence receipt

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.

receipt.json

{
  "claim": "candidate lowers buffer-rotated triad_f32 kernel latency versus baseline",
  "source": {
    "repository": "<repository URL>",
    "commit": "<full commit SHA>",
    "file_sha256": "<triad_bench.cu SHA-256>",
    "compiler": "nvcc <version>",
    "flags": ["-O3", "-lineinfo", "-std=c++17"]
  },
  "hardware": {
    "gpu_uuid": "<GPU UUID>",
    "gpu_name": "<exact model>",
    "compute_capability": "<major.minor>",
    "driver": "<driver version>",
    "cuda_runtime": "<runtime version>",
    "power_limit_w": "<watts>",
    "clock_policy": "observed or locked",
    "isolation": "<exclusive host, MIG slice, or competing-work policy>"
  },
  "workload": {
    "operation": "y[i] = fmaf(1.25f, x[i], z[i])",
    "dtype": "float32",
    "n": 1000003,
    "threads_per_block": 256,
    "cache_regime": "buffer-rotated",
    "rotation_slots": "<measured run value>"
  },
  "correctness": {
    "reference": "CPU double FMA, one cast to float",
    "atol": 0.000002,
    "rtol": 0.00002,
    "failures": 0,
    "memcheck_errors": 0
  },
  "measurement": {
    "timer": "CUDA events in kernel stream",
    "warmup": "<count and stability rule>",
    "samples": 200,
    "rounds": "<count>",
    "order": "ABBA within each round",
    "raw_samples": ["baseline-rotate.csv", "candidate-rotate.csv"]
  },
  "artifacts": {
    "systems": "triad_systems.nsys-rep",
    "compute": "triad_compute.ncu-rep",
    "sanitizer": "memcheck-fixed.txt",
    "gpu_state_before": "gpu-state-before.csv",
    "gpu_state_after": "gpu-state-after.csv"
  },
  "result": {
    "baseline_median_us": "<value>",
    "candidate_median_us": "<value>",
    "round_speedup_median": "<value>",
    "round_speedup_interval": "<method and bounds>",
    "end_to_end_effect": "<measured separately or explicitly out of scope>"
  }
}

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 small Python fixture below applies a release gate to one receipt. Its numbers are labeled fixture data; they aren't GPU measurements.

check-receipt-gate.py

from statistics import median

receipt_fixture = {
    "baseline_us": [102.0, 101.0, 104.0, 100.0, 103.0],
    "candidate_us": [82.0, 81.0, 83.0, 80.0, 82.0],
    "correctness_failures": 0,
    "sanitizer_errors": 0,
    "same_workload": True,
    "same_environment": True,
    "raw_samples_saved": True,
}

baseline = median(receipt_fixture["baseline_us"])
candidate = median(receipt_fixture["candidate_us"])
speedup = baseline / candidate
gates = [
    receipt_fixture["correctness_failures"] == 0,
    receipt_fixture["sanitizer_errors"] == 0,
    receipt_fixture["same_workload"],
    receipt_fixture["same_environment"],
    receipt_fixture["raw_samples_saved"],
]

print("fixture baseline median:", baseline, "us")
print("fixture candidate median:", candidate, "us")
print(f"fixture speedup: {speedup:.2f}x")
print("claim ready:", all(gates))

Output

fixture baseline median: 102.0 us
fixture candidate median: 82.0 us
fixture speedup: 1.24x
claim ready: True

Flip any gate to False. The speedup remains numerically the same, but the claim becomes unpublishable. That's intended behavior.

A defensible optimization loop

Keep the loop narrow:

State claim. Name workload, baseline, scope, correctness, and environment.
Prove semantics. Run reference comparison across adversarial shapes and values.
Check execution. Require relevant Compute Sanitizer tools to pass.
Find scope. Use Nsight Systems to locate delay on application critical path.
Explain kernel. Use Nsight Compute and Roofline only for confirmed hotspot.
Change one mechanism. Target measured traffic, launch cost, occupancy constraint, or math pipeline.
Rerun clean. Warm, control cache regime, record clocks and power, preserve raw samples.
Archive receipt. Link correctness, reports, environment, and result to exact source revision.

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.

Next Step

Continue to GPU Kernel Performance Engineering Lab

You can now produce a defensible baseline and trace one bottleneck without mistaking profiler output for proof. The next lab holds launch geometry fixed while changing memory path and kernel structure one mechanism at a time.

PreviousAccelerator Architecture Field Guide

Share this article

X Facebook LinkedIn Bluesky Reddit Hacker News Email

References

CUDA Programming Guide.

NVIDIA · 2026

https://docs.nvidia.com/cuda/cuda-programming-guide/

torch.allclose

PyTorch Contributors · 2026 · Official documentation

https://docs.pytorch.org/docs/stable/generated/torch.allclose

Compute Sanitizer

NVIDIA · 2026 · Official documentation

https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/

Nsight Systems User Guide

NVIDIA · 2026 · Official documentation

https://docs.nvidia.com/nsight-systems/UserGuide/

Nsight Compute Profiling Guide

NVIDIA · 2026 · Official documentation

https://docs.nvidia.com/nsight-compute/ProfilingGuide/

Roofline: An Insightful Visual Performance Model for Multicore Architectures

Williams, S., Waterman, A., & Patterson, D. · 2009

https://doi.org/10.1145/1498765.1498785

CUDA semantics.

PyTorch Contributors · 2026

https://docs.pytorch.org/docs/2.13/notes/cuda.html

PyTorch Benchmark

PyTorch Contributors · 2026 · Official tutorial

https://docs.pytorch.org/tutorials/recipes/recipes/benchmark.html

nvidia-smi documentation

NVIDIA · 2026

https://docs.nvidia.com/deploy/nvidia-smi/index.html

Discussion

Questions and insights from fellow learners.

Discussion loads when you reach this section.

Correct GPU Profiling and Benchmarking

Turn “faster” into a testable claim

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?

Prove output before studying speed

Choose instrument by question

Nsight Systems finds scope

Nsight Compute explains one kernel

Sanitizers run outside timing

Read Roofline and SpeedOfLight without guessing

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?

Design benchmark around real regime

Synchronize boundaries, not every operation

Warm until mechanism is steady

Name cache policy

Record clocks, power, and thermal state

Preserve distribution

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?

Run the triad investigation

Diagnose contradictory evidence

Preserve evidence receipt

A defensible optimization loop

Mastery check

Evaluation rubric

Common pitfalls

Mastery Check

Discussion

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?

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?

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?

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?

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?

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?

Mastery Check

Discussion

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?

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?