Measure your GPU's roofline

Three numbers, and then one prediction you check against them.

1. Achieved π — peak TFLOP/s, from a GEMM sweep that walks M from 1 to 16,384 against both square shapes and Llama-3-8B's real projection shapes.
2. Achieved β — peak GB/s, from two streaming kernels: a copy (one read, one write) and a triad (two reads, one write).
3. Your ridge point I* = π/β, the arithmetic intensity at which the bottleneck flips from HBM to tensor cores. Derivation in §0.4.

Then the check: the script places the book's three reference workloads on your roofline and prints T*, the token count at which a square d×d projection becomes compute-bound. On H100 spec-sheet numbers with d = 4096 that is 345 tokens. Your measured π will be lower than the datasheet's and your measured β will be lower too — but not by the same factor, so T* moves, and it moves in a direction the datasheet cannot tell you.

The skinny shapes are not invented. They are vLLM's own record of Llama-3.1-8B's projections, as [K, N] pairs with the tensor-parallel split dimension:

    "meta-llama/Llama-3.1-8B-Instruct": [
        ([4096, 6144], 1),
        ([4096, 4096], 0),
        ([4096, 28672], 1),
        ([14336, 4096], 0),
    ],

$ python3 run.py                                    # both halves, autodetect the card
$ python3 run.py --mode gemm --dtype float16        # fp16 tensor cores instead of bf16
$ python3 run.py --mode bandwidth --buffer-mib 4096 # bigger buffer, further from L2
$ python3 run.py --peak-tflops 989.4 --peak-gbps 3350 --csv roofline.csv
$ python3 run.py --help                             # every knob

Predict first, in this order, and write the answers down before you run anything.

1. What fraction of your card's published bf16 peak will the largest square GEMM reach?
2. What fraction of published HBM bandwidth will a copy reach? Will the triad be higher or lower?
3. Will your measured ridge point be higher or lower than the datasheet's?

The third is the interesting one, and most people get it wrong. It asks which of the two peaks is harder to reach in practice, and the answer is not obvious from either datasheet.

Timing, and the one mistake that invalidates everything

CUDA launches are asynchronous. A perf_counter wrapped around a kernel launch measures how fast Python enqueues work, which for a small GEMM is routinely faster than the work itself. run.py records CUDA events on the stream and reads them after a synchronize. Both engines' own harnesses do the equivalent — SGLang's single-batch harness calls model_runner.synchronize() on both sides of every timed region (§0.4 quotes it), and vLLM's kernel benchmarks delegate to Triton's do_bench, which does the same thing and returns a median rather than a mean.

The FLOP and byte accounting

FLOPs are 2·M·N·K: one multiply and one add per MAC. That is exactly vLLM's own accounting in the sweep this lab's GEMM half is modelled on — note that its x_vals straddle the T* ≈ 345 crossing, which is why the printed TFLOP/s column is the roofline's y-axis, sampled:

@triton.testing.perf_report(
    triton.testing.Benchmark(
        x_names=["batch_size"],
        x_vals=[1, 16, 64, 128, 256, 512, 1024, 2048, 4096, 8192, 16384],
        x_log=False,
        line_arg="provider",
        line_vals=_enabled,
        line_names=_enabled,
        ylabel="TFLOP/s (larger is better)",
# ...
    to_tflops = lambda t_ms: (2 * M * N * K) * 1e-12 / (t_ms * 1e-3)
    return to_tflops(ms), to_tflops(max_ms), to_tflops(min_ms)

Bytes are compulsory traffic — read A, read B, write C, each counted once. The same bytes-over-latency idiom appears in vLLM's cache-gather benchmark, which is the closest thing in either tree to a bandwidth probe:

    latency_ms = triton.testing.do_bench(
        run, warmup=warmup_ms, rep=rep_ms, return_mode="median"
    )
    bandwidth_gbps = bytes_moved / latency_ms / 1e6
    lengths = ",".join(str(seq_len) for seq_len in seq_lens)
    print(
        f"{variant:15s} {name:10s} batch={len(seq_lens):2d} "
        f"total={sum(seq_lens):7d} latency={latency_ms * 1e3:9.2f} us "
        f"bandwidth={bandwidth_gbps:8.1f} GB/s lengths=[{lengths}]"

The engines will tell you the same two numbers at runtime

Once you have π and β from bare kernels, both engines can report the versions they achieve while actually serving. vLLM's is a single log line, gated behind --enable-mfu-metrics:

    enable_mfu_metrics: bool = False
    """Enable Model FLOPs Utilization (MFU) metrics."""

        else:
            avg_tflops_per_gpu = self.total_num_flops_per_gpu / delta_time / 1e12
            avg_gbps_per_gpu = (
                (self.total_read_bytes_per_gpu + self.total_write_bytes_per_gpu)
                / delta_time
                / 1e9
            )

        log_fn(
            "%sMFU: %.1f TF/s/GPU %.1f GB/s/GPU",
            log_prefix,
            avg_tflops_per_gpu,
            avg_gbps_per_gpu,
        )

SGLang appends its equivalent to the decode-batch line, behind a flag of the same name, and splits read from write bandwidth — which is the more useful split, because decode is almost entirely reads:

        if self.enable_mfu_metrics and gap_latency > 0:
            flops_per_s = self._mfu_log_flops / gap_latency
            read_bytes_per_s = self._mfu_log_read_bytes / gap_latency
            write_bytes_per_s = self._mfu_log_write_bytes / gap_latency
            tflops_per_s = flops_per_s / 1e12
            read_gb_per_s = read_bytes_per_s / 1e9
            write_gb_per_s = write_bytes_per_s / 1e9
            msg += (
                f", est. decode TFLOPS/s (per GPU): {tflops_per_s:.2f}, "
                f"est. read BW (GB/s per GPU): {read_gb_per_s:.2f}, "
                f"est. write BW (GB/s per GPU): {write_gb_per_s:.2f}"
            )

Both are estimated — the engines count FLOPs and bytes analytically from the batch shape and divide by elapsed time; neither reads a hardware counter. So they are the same class of number as this lab's, computed on a real workload instead of a synthetic one. Comparing the two is the most direct way to find out how much of your card a serving step actually uses.

$ vllm serve meta-llama/Meta-Llama-3-8B-Instruct --enable-mfu-metrics 2>&1 | grep MFU
$ python3 -m sglang.launch_server --model-path meta-llama/Meta-Llama-3-8B-Instruct \
    --enable-mfu-metrics 2>&1 | grep "est. decode TFLOPS"

Neither measured number will match its datasheet, and the two miss by different amounts. That asymmetry is the result.

Then substitute. Every KPI in Part 0 and Part 1 was computed against 989.4 and 3.35 TB/s. The Llama-3-8B decode floor of 4.48 ms per step, the 223 tok/s batch-1 ceiling, the T* = 345 crossing, the 295 FLOP/byte ridge — all of them scale with the two numbers you just measured, and all of them get worse. That is not a flaw in the derivations; it is what "spec-sheet peak" means.