LLM inference optimization

From first principles to interview depth. Every number is either computed with a script, measured on our stg RTX PRO 6000 Blackwell, or explicitly marked as a vendor spec.

1. Serving metrics & SLO vocabulary

The central tension

Bigger batches raise throughput (weights are read once and amortized over more tokens) but raise per-request latency. Every serving decision is a move along that curve. The job is to find the highest-batch operating point that still meets the p99 SLO, which is why the right cost metric is cost at SLO.

Copy# computed
e2e_ms  = 300 + (200 - 1) * 25                 # TTFT 300 ms, TPOT 25 ms, 200 tokens -> 5275 ms
usd_per_M = 3.0 / (5000 * 3600) * 1e6           # $3/hr GPU, 5000 tok/s -> $0.167 per 1M tokens
p_any_slow = 1 - 0.99 ** 100                    # 100 independent calls, each p99 -> 63.4% see a slow one

2. Roofline for LLMs: prefill vs decode

Arithmetic intensity (AI) = FLOPs performed per byte moved from memory. A GPU has a compute peak P (FLOP/s) and a bandwidth B (bytes/s). Attainable speed = min(P, B x AI). The ridge point P/B is the AI where you switch from memory-bound to compute-bound.

Why decode is memory-bound

Generating one token for one sequence multiplies a [1 x d] activation by every weight matrix. Each weight is read once from HBM/GDDR and used for ~2 FLOPs (one multiply, one add). AI is about 1-2 FLOP/byte, far below any ridge point (hundreds). Adding more sequences to the batch reuses the same weight read: a batch of B sequences does B times the FLOPs for the same bytes, so for an fp16 weight GEMM AI is about B FLOP/byte.

Prefill processes the whole prompt as a [T x d] matrix, so AI is about T and it sits at the compute roof.

The floor formula

time per decode step >= bytes_read_per_step / memory_bandwidth. For a single sequence the bytes are essentially the weight bytes (plus KV bytes). It is a floor: real systems land at 60-85% of peak bandwidth, plus launch and sampling overhead.

All table values are computed from the spec bandwidths. The L4 is 6x slower on every row because its bandwidth is 6x lower, which matches the earlier L4-vs-RTX 6000 result for being memory-bound.

Lab 2a: decode-like GEMM, time vs batch (measured)

Weight matrix 8192 x 8192 fp16 (134 MB, rotated through 4 copies so L2 cannot serve repeats). x @ W.T with x of shape [B, 8192]. This is exactly one linear layer of a decode step at batch B.

Copyimport torch
def bench(f, n=50):
    for _ in range(5): f()
    torch.cuda.synchronize()
    s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
    s.record()
    for _ in range(n): f()
    e.record(); torch.cuda.synchronize()
    return s.elapsed_time(e) / n * 1e3  # us
torch.manual_seed(0)
N = K = 8192
# 4 distinct weight copies cycled so L2 (~100+ MB class) cannot serve repeats
Ws = [torch.randn(N, K, device="cuda", dtype=torch.float16) for _ in range(4)]
print(f"weight bytes per matrix: {N*K*2/1e6:.1f} MB (fp16)")
print(f"{'B':>5} {'time_us':>9} {'GB/s':>8} {'TFLOPS':>8}")
for B in [1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096]:
    x = torch.randn(B, K, device="cuda", dtype=torch.float16)
    i = [0]
    def f():
        i[0] = (i[0] + 1) % 4
        return x @ Ws[i[0]].T
    t = bench(f)
    print(f"{B:5d} {t:9.1f} {N*K*2/t/1e3:8.0f} {2*B*N*K/t/1e6:8.1f}")

measured on RTX PRO 6000 Blackwell, stg

weight bytes per matrix: 134.2 MB (fp16)
    B   time_us     GB/s   TFLOPS
    1      97.9     1371      1.4
    2     100.3     1339      2.7
    4     100.2     1339      5.4
    8      99.7     1346     10.8
   16      99.5     1349     21.6
   32     103.3     1300     41.6
   64      97.8     1372     87.8
  128     114.3     1174    150.3
  256     123.7     1085    277.7
  512     206.4      650    332.9
 1024     359.2      374    382.6
 2048     694.4      193    395.8
 4096    1437.3       93    382.5

Prefill TTFT estimate

FLOPs for a forward pass of T tokens is about 2 x params x T (ignoring attention). 8B model, 2000-token prompt: 2 x 8e9 x 2000 = 3.2e13 FLOP; at 400 TFLOPS = 80 ms, at a more realistic 200 TFLOPS = 160 ms (computed). For a 70B model the same prompt is 2.8e14 FLOP, i.e. 700 ms at 400 TFLOPS on one GPU. That single fact drives much of section 13's design question.

3. KV cache

Attention for the new token needs the keys and values of every previous token. Recomputing them each step would be quadratic work, so we store them: the KV cache.

Copybytes_per_token = 2 * n_layers * n_kv_heads * head_dim * bytes_per_elem   # 2 = K and V
total_bytes     = bytes_per_token * tokens_in_cache   # summed over all live sequences

One 8k-token sequence on the 70B-like model = 2.5 GiB of KV; on the 8B-like model, 1.0 GiB (computed).

Lab 3: the formula vs the allocator (measured)

Copyimport torch
layers, kvh, hd, bytes_, seqs, toks = 32, 8, 128, 2, 4, 4096
per_tok = 2 * layers * kvh * hd * bytes_
pred = per_tok * seqs * toks
torch.cuda.synchronize(); base = torch.cuda.memory_allocated()
kv = torch.empty(layers, 2, seqs, toks, kvh, hd, dtype=torch.float16, device="cuda")
meas = torch.cuda.memory_allocated() - base
print(f"formula: 2*{layers}*{kvh}*{hd}*{bytes_} = {per_tok} B/token = {per_tok/1024:.0f} KiB/token")
print(f"predicted for {seqs} seqs x {toks} tok: {pred} B = {pred/2**30:.3f} GiB")
print(f"measured memory_allocated delta   : {meas} B = {meas/2**30:.3f} GiB  match={pred==meas}")

measured on RTX PRO 6000 Blackwell, stg

formula: 2*32*8*128*2 = 131072 B/token = 128 KiB/token
predicted for 4 seqs x 4096 tok: 2147483648 B = 2.000 GiB
measured memory_allocated delta   : 2147483648 B = 2.000 GiB  match=True

GQA / MQA

Multi-head attention keeps K,V per query head. MQA shares one K/V head across all query heads; GQA shares one per group (e.g. 32 query heads over 8 KV heads). Quality loss is small, and KV size and decode KV bandwidth drop by the group factor (4x in the table above). Since decode reads the whole KV cache every step, GQA speeds up decode as well as shrinking memory.

Why vLLM preallocates

• Memory is the batch-size limit: concurrent sequences = free memory / KV per sequence. Whatever is left after weights and activations is converted into KV blocks at startup (gpu_memory_utilization, default 0.9 in vLLM as I recall). Hence "vLLM fills the card".
• Dynamic cudaMalloc/free of large variable-size tensors fragments memory and stalls the stream (cudaMalloc can synchronize). One big pool avoids both.
• Fixed-size blocks make allocation O(1) and fragmentation bounded. That is PagedAttention.

PagedAttention

Naively each sequence reserves a contiguous slab for its maximum length: most is unused (internal fragmentation) and slabs of different sizes leave holes (external fragmentation). The vLLM paper reported large waste in earlier systems (I recall 60-80%; treat as from memory). PagedAttention splits the cache into fixed blocks (typically 16 tokens) and gives each sequence a block table mapping logical position to physical block, like virtual memory pages. The attention kernel indexes through the table. Consequences: waste is at most one partial block per sequence; blocks can be shared (prefix caching, beam search, parallel sampling) with copy-on-write; sequences can grow without moving.

4. Continuous batching, chunked prefill, scheduling

Static batching: collect N requests, run them together until the longest finishes. Short requests wait and the GPU idles on finished rows. Continuous (in-flight, iteration-level) batching (Orca): the scheduler runs every step; finished sequences leave and waiting ones join at the next iteration boundary. Throughput rises several-fold and queueing delay falls, because there is no head-of-line blocking behind the longest sequence.

The prefill/decode interference problem

A step that includes a 8000-token prefill is compute-heavy and long (hundreds of ms). Every decoding user in that batch sees a stalled token: an ITL spike that wrecks p99 TPOT.

Chunked prefill: cap the tokens per step (max_num_batched_tokens), fill the budget first with all decode tokens (1 each), then with a chunk of pending prefill. A long prompt is split across several steps, each with bounded length. Trade-off: it raises that request's TTFT slightly (more steps, each also carrying decode work) in exchange for flat ITL for everyone else. A bigger budget favours TTFT/throughput; a smaller one favours ITL. Recent vLLM versions enable it by default (from memory; check your version's defaults).

Example (computed): budget 2048, 100 running decodes leaves 1948 tokens/step for prefill, so an 8000-token prompt needs ceil(8000/1948) = 5 steps.

Scheduling policy and preemption

• Default is FCFS with a running set and a waiting queue. Admission needs free KV blocks.
• When KV runs out mid-flight (sequences grew), the scheduler preempts: it evicts a sequence's blocks and either recomputes it later (re-prefill; cheap for short contexts, and it can reuse prefix cache) or swaps blocks to CPU memory (costs PCIe bandwidth, section 11 shows ~55 GB/s pinned). Recompute is the common default in V1 as I recall.
• Preemption shows up as unexplained p99 spikes and a rising "num preempted" metric. Fixes: more KV room (quantize KV, smaller max_model_len, higher memory utilization), a lower max_num_seqs, or admission control upstream.
• Priority / SLO-aware scheduling, and separating long-prompt from short-prompt traffic, are the next levers.

5. Attention kernels: FlashAttention, FlashDecoding

Naive attention: S = QK^T (N x N), P = softmax(S), O = PV. The N x N matrix is written to HBM, read back for softmax, written, read again for PV. Cost is dominated by memory traffic, not FLOPs, and memory footprint is O(N^2).

FlashAttention (IO-aware): split Q into row blocks and K,V into column blocks that fit in on-chip SRAM/shared memory. For each Q block, loop over K/V blocks, compute the score tile on-chip, update the output, and never write S to HBM. Same FLOPs (it is exact attention, not an approximation), far fewer HBM bytes, O(N) extra memory.

The trick that makes tiling possible is online softmax. Keep per-row running max m and running sum l. For a new tile with scores s:

Copym_new = max(m, rowmax(s))
p     = exp(s - m_new)
l     = l * exp(m - m_new) + rowsum(p)
O     = O * exp(m - m_new) + p @ V_tile
m     = m_new
# at the end: O / l   (identical to softmax(S) @ V)

Lab 5: naive vs online-softmax tiling vs fused SDPA (measured)

Copyimport torch, math, torch.nn.functional as F
torch.manual_seed(0)
H, N, D = 16, 8192, 128
q, k, v = (torch.randn(1, H, N, D, device="cuda", dtype=torch.float16) for _ in range(3))

def naive(q, k, v):  # materializes the full N x N score matrix per head in HBM
    s = (q @ k.transpose(-1, -2)) / math.sqrt(D)
    return torch.softmax(s.float(), -1).half() @ v

def tiled(q, k, v, BK=1024):  # online softmax over K/V tiles; never builds N x N
    scale = 1 / math.sqrt(D)
    m = torch.full(q.shape[:-1] + (1,), -float("inf"), device="cuda", dtype=torch.float32)
    l = torch.zeros_like(m); o = torch.zeros(q.shape, device="cuda", dtype=torch.float32)
    for j in range(0, N, BK):
        s = (q @ k[..., j:j+BK, :].transpose(-1, -2)).float() * scale
        m_new = torch.maximum(m, s.amax(-1, keepdim=True))
        p = torch.exp(s - m_new); corr = torch.exp(m - m_new)
        l = l * corr + p.sum(-1, keepdim=True)
        o = o * corr + p.half().float() @ v[..., j:j+BK, :].float()
        m = m_new
    return (o / l).half()

def run(f):
    torch.cuda.reset_peak_memory_stats(); torch.cuda.synchronize()
    base = torch.cuda.memory_allocated()
    for _ in range(2): f()
    s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
    s.record()
    for _ in range(5): out = f()
    e.record(); torch.cuda.synchronize()
    return out, s.elapsed_time(e) / 5, (torch.cuda.max_memory_allocated() - base) / 2**20

ref, t_sdpa, m_sdpa = run(lambda: F.scaled_dot_product_attention(q, k, v))
o1, t_n, m_n = run(lambda: naive(q, k, v))
o2, t_t, m_t = run(lambda: tiled(q, k, v))
print(f"H={H} N={N} D={D} fp16 (prefill-shaped, non-causal)")
print(f"{'impl':28s} {'ms':>8} {'peak extra MiB':>15} {'max|err| vs SDPA':>18}")
print(f"{'SDPA (fused flash kernel)':28s} {t_sdpa:8.2f} {m_sdpa:15.0f} {'-':>18}")
print(f"{'naive (N x N in HBM)':28s} {t_n:8.2f} {m_n:15.0f} {(o1-ref).abs().max().item():18.4f}")
print(f"{'python tiled online-softmax':28s} {t_t:8.2f} {m_t:15.0f} {(o2-ref).abs().max().item():18.4f}")
print(f"score matrix alone would be {H*N*N*2/2**20:.0f} MiB in fp16")

measured on RTX PRO 6000 Blackwell, stg

H=16 N=8192 D=128 fp16 (prefill-shaped, non-causal)
impl                               ms  peak extra MiB   max|err| vs SDPA
SDPA (fused flash kernel)        1.82              65                  -
naive (N x N in HBM)            20.62           10280             0.0003
python tiled online-softmax     44.90            2146             0.0002
score matrix alone would be 2048 MiB in fp16

FlashDecoding

FlashAttention parallelizes over (batch, head, Q blocks). In decode there is one query per sequence, so batch x heads may be only 8-64 blocks on a GPU with ~100+ SMs, while the KV length is thousands of tokens. FlashDecoding additionally splits the KV sequence across thread blocks; each computes a partial (m, l, O), and a small reduction combines them using the same rescaling identity. It fills the GPU for long-context, small-batch decode. (Modern libraries like FlashInfer and vLLM's paged-attention kernels use split-KV in this spirit; implementation details differ and I have not read the current ones.)

6. Quantization

Lab 6: 8-bit vs fp16 on the same GEMM shape (measured)

Same 8192 x 8192 weight, 4 rotating copies. int8 via torch._int_mm, fp8 (e4m3) via torch._scaled_mm, random data (timing only, not accuracy). _int_mm requires B > 16, hence the error at B=16.

Copyimport torch
def bench(f, n=50):
    for _ in range(5): f()
    torch.cuda.synchronize()
    s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
    s.record()
    for _ in range(n): f()
    e.record(); torch.cuda.synchronize()
    return s.elapsed_time(e) / n * 1e3
N = K = 8192
for B in [16, 64, 256, 2048]:
    x16 = torch.randn(B, K, device="cuda", dtype=torch.float16)
    W16 = [torch.randn(N, K, device="cuda", dtype=torch.float16) for _ in range(4)]
    c = [0]
    def f16():
        c[0] = (c[0]+1) % 4; return x16 @ W16[c[0]].T
    t16 = bench(f16)
    row = f"B={B:5d} fp16 {t16:8.1f}us"
    try:
        x8 = torch.randint(-127, 127, (B, K), device="cuda", dtype=torch.int8)
        W8 = [torch.randint(-127, 127, (N, K), device="cuda", dtype=torch.int8) for _ in range(4)]
        def f8():
            c[0] = (c[0]+1) % 4; return torch._int_mm(x8, W8[c[0]].T)
        t8 = bench(f8); row += f" | int8 W8A8 {t8:8.1f}us ({t16/t8:.2f}x)"
    except Exception as ex: row += f" | int8 err {str(ex)[:60]}"
    try:
        xf = x8.float().to(torch.float8_e4m3fn)
        Wf = [W.float().to(torch.float8_e4m3fn) for W in W8]
        sa = torch.tensor(1.0, device="cuda"); sb = torch.tensor(1.0, device="cuda")
        def ff():
            c[0] = (c[0]+1) % 4; return torch._scaled_mm(xf, Wf[c[0]].T, scale_a=sa, scale_b=sb, out_dtype=torch.float16)
        tf = bench(ff); row += f" | fp8 W8A8 {tf:8.1f}us ({t16/tf:.2f}x)"
    except Exception as ex: row += f" | fp8 err {str(ex)[:80]}"
    print(row)
    del W16

measured on RTX PRO 6000 Blackwell, stg

B=   16 fp16     99.9us | int8 err self.size(0) needs to be greater than 16, but got 16 | fp8 W8A8     47.8us (2.09x)
B=   64 fp16     97.9us | int8 W8A8     47.5us (2.06x) | fp8 W8A8     48.8us (2.00x)
B=  256 fp16    123.5us | int8 W8A8     66.3us (1.86x) | fp8 W8A8     66.6us (1.86x)
B= 2048 fp16    727.5us | int8 W8A8    375.6us (1.94x) | fp8 W8A8    363.2us (2.00x)

Rules of thumb

• Latency-critical small-batch decode: W4A16 (AWQ/GPTQ) gives the biggest gain per accuracy point lost, which matches why INT4 AWQ gemma-3-1b is a sensible production choice.
• High-batch throughput serving: FP8 W8A8 (plus FP8 KV) because it helps in both regimes.
• Always re-certify accuracy on your task: quantization damage is task-dependent, and small models (1B) have less redundancy than 70B.
• AWQ protects salient weight channels using activation statistics; GPTQ minimizes layer-wise reconstruction error with second-order information. Both need a small calibration set.

7. Speculative decoding

Decode is memory-bound, so verifying k extra tokens costs almost the same as generating one (the weights are read once either way; lab 2a shows time flat up to B=64). Idea: a cheap draft proposes k tokens; the big target model scores all k+1 positions in one forward pass; accept the longest prefix that matches what the target would have sampled (rejection sampling keeps the output distribution exactly the target's). Each step yields at least 1 token and up to k+1.

Variants

• Separate draft model (e.g. 1B drafting for 70B, same tokenizer). Simple; draft needs its own KV and compute.
• Medusa: extra heads on the target's last hidden state each predict a future token; verified as a tree of candidates. No second model, but heads need training.
• EAGLE: a small autoregressive head that predicts the next feature (hidden state) using the target's hidden states, giving higher acceptance than token-only drafting; also tree-verified. (From memory; read the papers before an interview.)
• n-gram / prompt-lookup: draft by copying spans from the prompt; free, great for summarization/RAG/code edits.

Acceptance math

If each drafted token is accepted independently with probability alpha, expected tokens per target step is (1 - alpha^(k+1)) / (1 - alpha). With draft cost c (relative to one target step), speedup is about E / (1 + k c) in the memory-bound regime.

Copyimport random
random.seed(1)
def sim(a, k, trials=200000):
    tot = 0
    for _ in range(trials):
        n = 0
        while n < k and random.random() < a: n += 1
        tot += n + 1  # +1: the target's own token (correction or bonus)
    return tot / trials
print(f"{'alpha':>5} {'k':>2} {'simulated':>10} {'formula (1-a^(k+1))/(1-a)':>27}")
for a in [0.5, 0.7, 0.8, 0.9]:
    for k in [3, 5]:
        print(f"{a:5.1f} {k:2d} {sim(a,k):10.3f} {(1-a**(k+1))/(1-a):27.3f}")
print()
c = 0.1  # draft step cost relative to one target step
print("speedup = E[tokens] / (1 + k*c), c=0.1, memory-bound regime (verify cost ~ one decode step)")
for a in [0.5, 0.7, 0.8, 0.9]:
    for k in [3, 5]:
        e = (1-a**(k+1))/(1-a); print(f"alpha={a} k={k}: {e/(1+k*c):.2f}x")

Monte-Carlo simulation, run locally (pure Python; not GPU)

alpha  k  simulated   formula (1-a^(k+1))/(1-a)
  0.5  3      1.873                       1.875
  0.5  5      1.966                       1.969
  0.7  3      2.534                       2.533
  0.7  5      2.948                       2.941
  0.8  3      2.949                       2.952
  0.8  5      3.692                       3.689
  0.9  3      3.436                       3.439
  0.9  5      4.691                       4.686

speedup = E[tokens] / (1 + k*c), c=0.1, memory-bound regime (verify cost ~ one decode step)
alpha=0.5 k=3: 1.44x
alpha=0.5 k=5: 1.31x
alpha=0.7 k=3: 1.95x
alpha=0.7 k=5: 1.96x
alpha=0.8 k=3: 2.27x
alpha=0.8 k=5: 2.46x
alpha=0.9 k=3: 2.65x
alpha=0.9 k=5: 3.12x

Simulation confirms the closed form. Note the diminishing returns: at alpha=0.5, going from k=3 to k=5 lowers the speedup (1.44x to 1.31x) because the extra draft cost buys almost no extra accepted tokens.

8. Parallelism & communication

Order of preference: fit on one GPU (quantize!) and replicate; else TP within a node; PP/EP across nodes. TP buys latency and memory at the cost of communication; DP buys throughput for free.

Communication cost, with numbers

An 80-layer model with TP does 2 x 80 = 160 all-reduces per token step. At batch 32 and hidden size 8192 in fp16 each message is only 32 x 8192 x 2 = 512 KiB (computed). Small messages are latency-bound, not bandwidth-bound: 160 collectives times a per-collective latency of tens of microseconds adds milliseconds per token, comparable to the entire compute floor of 2-5 ms. That is why TP wants NVLink-class links (high bandwidth and low latency, with in-network/one-shot all-reduce tricks) and degrades on PCIe, where bandwidth and latency are both worse and traffic may traverse the CPU root complex.

MoE serving

• All experts' weights must be resident (memory), but each token only activates top-k (e.g. 2 of 8, or 8 of 256) so FLOPs per token are far lower than total parameters suggest.
• At batch 1 only the selected experts' bytes are read, so decode is cheap. As batch grows, nearly every expert is hit each step, so bytes read per step approach the full model: the memory-bound advantage fades with batch size.
• Expert parallelism spreads experts across GPUs; each layer needs an all-to-all of token activations. Hot experts cause stragglers (the step is gated by the busiest GPU). Techniques: expert replication, load-balanced routing, overlapping communication with compute.

9. Prefix caching & disaggregated prefill/decode

Prefix caching

Many requests share a long prefix (system prompt, few-shot examples, a document, chat history). KV for identical prefix tokens is identical, so keep it. vLLM hashes each full block of tokens (chained with the previous block's hash) to find reusable physical blocks; unreferenced blocks are evicted LRU. Savings: prefill compute for the prefix (= TTFT and GPU time).

• Only full blocks are reused: a 2000-token prompt with an 1800-token shared prefix and 16-token blocks reuses 112 blocks = 1792 tokens and recomputes 208 (computed).
• Hit rate depends on prompt layout (put stable content first, variable content last) and on routing: with N replicas, requests with the same prefix must reach the same replica (prefix-aware / consistent-hash routing) or you divide your hit rate by roughly N.
• The cache lives in the same pool as live KV; heavy caching reduces batch capacity.
• This is the serving analogue of your caching experience: key design, hit rate, eviction, and cache-aware load balancing.

Disaggregated prefill/decode

Prefill is compute-bound and bursty; decode is memory-bound and steady. Co-located, they interfere (chunked prefill mitigates but does not remove it) and you must pick one parallelism and one hardware type for both. Disaggregation runs prefill on one pool and decode on another and ships the KV cache between them.

• Appeal: independent scaling and tuning (prefill pool sized for TTFT, decode pool for TPOT); hardware matched to the bottleneck (compute-rich for prefill, bandwidth-rich for decode); decode ITL no longer disturbed by prefills.
• Cost: KV transfer. For an 8B-like model and a 2000-token prompt the KV is 131072 x 2000 = 262 MB. Over a 25 Gbit/s link that is 83.9 ms; at 400 Gbit/s, 5.2 ms; over a 55 GB/s PCIe-class path, 4.8 ms (computed). Needs RDMA/NVLink-class transfer, layer-wise streaming to overlap with compute, and a more complex scheduler and failure model.
• Pays off at larger scale, long prompts, strict dual SLOs. For small models, short prompts or few GPUs it is usually not worth it (my judgment, not a measured result).

10. Kernels & compilers

10a. CUDA graphs: launch overhead (measured)

Each kernel launch costs the CPU several microseconds. Decode runs hundreds of small kernels per step, so launch overhead can exceed GPU time. A CUDA graph records a kernel sequence once and replays it with a single launch.

Copyimport torch
x = torch.randn(1024, device="cuda")
M = 300
def body(x):
    for _ in range(M): x = x * 1.0001 + 0.0001
    return x
def bench(f, n=20):
    for _ in range(3): f()
    torch.cuda.synchronize()
    s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
    s.record()
    for _ in range(n): f()
    e.record(); torch.cuda.synchronize()
    return s.elapsed_time(e) / n
te = bench(lambda: body(x))
g = torch.cuda.CUDAGraph()
st = torch.cuda.Stream()
with torch.cuda.stream(st):
    for _ in range(3): body(x)
torch.cuda.synchronize()
with torch.cuda.graph(g): y = body(x)
tg = bench(lambda: g.replay())
print(f"{2*M} tiny kernels ({M} iterations of mul+add on 1024 floats)")
print(f"eager      : {te:.3f} ms per call = {te/(2*M)*1e3:.2f} us per kernel")
print(f"cuda graph : {tg:.3f} ms per call = {tg/(2*M)*1e3:.2f} us per kernel")
print(f"speedup    : {te/tg:.1f}x")

measured on RTX PRO 6000 Blackwell, stg

600 tiny kernels (300 iterations of mul+add on 1024 floats)
eager      : 3.633 ms per call = 6.05 us per kernel
cuda graph : 0.612 ms per call = 1.02 us per kernel
speedup    : 5.9x

These are deliberately tiny kernels (1024 floats), so this is the worst case for launch overhead: ~6 us per kernel eager is CPU-launch-bound, ~1 us per kernel in the graph is close to the GPU's own per-kernel floor. Real decode kernels are larger, so the benefit is smaller but still large for small models, which is exactly where the 1B-class models live. Constraints: static shapes and static addresses, so vLLM captures graphs for a set of batch sizes and pads up; prefill (variable shapes) typically runs eagerly. Nsight Systems shows the symptom as gaps between kernels on the GPU timeline.

10b. Fusion: Triton and torch.compile (measured)

y = x*2 + 1 in eager PyTorch is two kernels: each reads and writes the full tensor, so 4 passes over memory instead of 2. A fused kernel does one read, one write.

Copyimport torch, triton, triton.language as tl
@triton.jit
def k(x_ptr, y_ptr, n, a, b, BLOCK: tl.constexpr):
    offs = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK)
    m = offs < n
    x = tl.load(x_ptr + offs, mask=m)
    tl.store(y_ptr + offs, x * a + b, mask=m)
def tri(x, y):
    n = x.numel(); k[(triton.cdiv(n, 1024),)](x, y, n, 2.0, 1.0, BLOCK=1024)
def bench(f, n=100):
    for _ in range(5): f()
    torch.cuda.synchronize()
    s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
    s.record()
    for _ in range(n): f()
    e.record(); torch.cuda.synchronize()
    return s.elapsed_time(e) / n * 1e3
n = 1 << 28  # 1 GiB fp32
x = torch.randn(n, device="cuda"); y = torch.empty_like(x)
tri(x, y); assert torch.allclose(y, x * 2 + 1)
byt = 2 * n * 4
t1 = bench(lambda: tri(x, y))
t2 = bench(lambda: x * 2.0 + 1.0)
cf = torch.compile(lambda x: x * 2.0 + 1.0)
cf(x); t4 = bench(lambda: cf(x))
t5 = bench(lambda: y.copy_(x))
for name, t, b in [("triton fused", t1, byt), ("torch eager x*2+1 (2 kernels)", t2, 4*n*4), ("torch.compile fused", t4, byt), ("plain copy (reference)", t5, byt)]:
    print(f"{name:32s} {t:8.0f} us  {b/t/1e3:7.0f} GB/s moved")

measured on RTX PRO 6000 Blackwell, stg (1 GiB fp32 tensor)

triton fused                         1470 us     1461 GB/s moved
torch eager x*2+1 (2 kernels)        2939 us     1461 GB/s moved
torch.compile fused                  1469 us     1462 GB/s moved
plain copy (reference)               1474 us     1457 GB/s moved

The tool landscape

TensorRT-LLM vs vLLM vs SGLang (honest view)

11. Multi-tenancy, LoRA, autoscaling & cold starts

LoRA multi-tenant serving

A LoRA adapter adds low-rank matrices A (r x d_in) and B (d_out x r) to selected linears: y = Wx + B(Ax). One base model in memory serves many tenants' adapters. Rank 16 on all linears of a Llama-3-8B-shaped model is 41.9M parameters, ~84 MB in fp16 (computed), 0.5% of the base, so thousands fit in CPU memory and dozens on GPU.

• The batch can contain requests for different adapters. Kernels such as BGMV/SGMV (punica/S-LoRA style) gather per-request adapter weights; the base GEMM is shared. Overhead is real but small at low rank.
• Scheduling: adapter hot-set on GPU, LRU to CPU, load on demand (an adapter miss is a mini cold start of tens of ms). Route requests to replicas that already have the adapter.
• Isolation concerns: fairness (one tenant's burst starves others), per-tenant rate limits and KV quotas, prefix-cache poisoning/side channels across tenants (cache keys should include tenant/adapter id).

Autoscaling LLM replicas

• CPU utilization is useless. GPU utilization is misleading (a GPU can be "100% busy" at 10% of its attainable throughput). Better signals: queue depth / waiting requests, KV-cache utilization, tokens in flight, TTFT/TPOT vs SLO.
• Scale-up is slow (minutes), so scale on leading indicators and keep headroom. Scale-down kills KV and prefix caches, so drain gracefully.
• Keep a minimum warm pool for latency-sensitive tenants; scale-to-zero trades idle-GPU cost against cold-start latency.

Cold start anatomy

1. Provision the GPU node (cloud API, minutes; mitigated with warm node pools).
2. Pull the container image (multi-GB; mitigated by lazy/streamed image layers and local caches).
3. Load weights: storage/network to host memory to GPU.
4. Initialize: CUDA context, allocate KV pool, torch.compile, capture CUDA graphs, warm kernels (often tens of seconds if not cached).

Lab 11: host to GPU copy (measured)

Copyimport torch, time
n = 2 * 2**30
for name, pin in [("pageable", False), ("pinned", True)]:
    h = torch.empty(n, dtype=torch.uint8, pin_memory=pin); h.fill_(1)
    d = torch.empty(n, dtype=torch.uint8, device="cuda")
    d.copy_(h, non_blocking=pin); torch.cuda.synchronize()
    t = time.time()
    for _ in range(3): d.copy_(h, non_blocking=pin)
    torch.cuda.synchronize(); dt = (time.time() - t) / 3
    print(f"{name:9s} H2D 2 GiB: {dt*1e3:7.1f} ms = {n/dt/1e9:5.1f} GB/s")

measured on RTX PRO 6000 Blackwell, stg

pageable  H2D 2 GiB:   206.2 ms =  10.4 GB/s
pinned    H2D 2 GiB:    38.7 ms =  55.5 GB/s

Last-hop PCIe is not the bottleneck: 16 GB over 55 GB/s is 0.29 s, versus 16 s if the source delivers 1 GB/s, and 35 s for 70 GB at 2 GB/s (computed). Pageable (non-pinned) memory is 5x slower, so loaders stage through pinned buffers. Cold-start work is therefore about the earlier hops, which is where the standard techniques act:

• Weight streaming / parallel fetch: many concurrent range reads from object storage, overlapped with host-to-GPU copies (load layer N while fetching N+1); formats like safetensors allow zero-copy mmap.
• Local/peer caches: node-local NVMe, a cluster cache, or peer-to-peer pulls so the object store is hit once per node, not per replica.
• Snapshotting: checkpoint an initialized process (including, in some systems, GPU memory state) and restore it instead of re-running init. Appealing because it skips compile/graph-capture/warmup, but GPU state checkpointing is intricate (driver, CUDA context and memory). I describe the idea; I have not verified any tool's maturity.
• Cache compiled artifacts (torch.compile/Inductor cache, Triton cache, CUDA graphs can't be serialized simply, so capture time is mostly paid at start).

12. Capacity planning & benchmarking methodology

Open vs closed loop

• Closed loop: N virtual users each send, wait for a response, then send again. When the server slows, the clients automatically slow too, so the offered load shrinks and the measured latency understates queueing (a form of coordinated omission). Good for finding max throughput at fixed concurrency.
• Open loop: requests arrive on a schedule (e.g. Poisson at rate lambda) regardless of server state. This matches real traffic, exposes queueing collapse, and is what you need to measure p99 at a given QPS. Idealized M/M/1 queueing: mean time in system is service time / (1 - utilization), so 50% load is 2x, 80% is 5x, 90% is 10x, 95% is 20x service time (computed). Real systems are not M/M/1 but the hockey-stick shape is real, which is why you cannot run GPUs at 95%.

Method checklist

1. Realistic length distributions: sample prompt and output lengths from production logs, not a fixed 128/128. Long-tail prompts drive p99 TTFT; output length drives KV occupancy.
2. Realistic content for caching: shared-prefix fraction determines prefix-cache hit rate; random prompts understate it, a single repeated prompt overstates it.
3. Warmup: discard the first requests (CUDA graph capture, compile, lazy loading, empty caches); measure steady state, long enough to include GC pauses and cache eviction (minutes, not seconds).
4. Sweep offered load (QPS) to produce a latency-vs-throughput curve; capacity = highest QPS where p99 meets SLO. Report goodput.
5. Look at percentiles, not means; collect enough samples (p99 needs hundreds+, p99.9 thousands). Report per-request-class percentiles (short vs long prompts).
6. Keep the load generator off the critical path: a saturated client (Python, tokenizer, TLS) makes the server look slow. Check client CPU.
7. Change one variable at a time; repeat runs; record versions, flags, GPU clocks.
8. Accuracy gate every perf experiment that touches numerics (quantization, spec decode, kernels).

Tied to your world

13. Interview prep

Open each answer only after you have said yours out loud. Answers lean on the numbers measured above.

How to talk about your experience