If you compared raw hidden states, matching would collapse into plain similarity — a token could only attend to tokens that look like itself. Separate Q and K projections let the model learn an asymmetric relation: hungry can ask for a subject without resembling one.

Dot products of d-dimensional vectors grow with d. Without the scaling, scores get large, softmax saturates to nearly one-hot, and gradients through it vanish. The divisor keeps the logits in a range where softmax stays soft and trainable.

“Why is it quadratic?” Because every query meets every key — S² pairs, each a d-dimensional dot product. That single fact is the origin of FlashAttention, sparse attention, linear attention, and most of the long-context literature.

Softmax attention cannot be reassociated because the softmax sits between Q·Kᵀ and V. Replace it with a feature map φ so scores become φ(q)·φ(k), and the computation reassociates: instead of (QKᵀ)V you compute Q(KᵀV). KᵀV is a fixed-size matrix independent of sequence length — that matrix is the recurrent state.

Retrieval-shaped work (copying an exact token, following a reference) is what a compressed state does worst. Interleaving a few full-attention layers among many linear ones recovers that ability while keeping most of the cost saving — which is why most published “linear” models are not purely linear.

Linear attention is not a strict improvement. It trades an exact, addressable memory for a lossy summary. If your workload is long-context retrieval, that trade is bad; if it is long-context summarisation, it is good.

Softmax normally needs the whole row before it can divide by the sum. Instead, carry a running max m and running sum ℓ. When a new tile arrives with a larger max, rescale the accumulated output by exp(m_old − m_new) before adding the new contribution. The correction is exact, so the final result is bit-for-bit what the naive version produces.

In the backward pass FlashAttention recomputes the attention tiles rather than storing them. That is more FLOPs but far less memory traffic — and since the kernel is bandwidth-bound, trading arithmetic for traffic is a win. It is the same reasoning as the forward pass, applied to gradients.

FlashAttention-1 introduced the tiled, fused, IO-aware formulation. FlashAttention-2 kept the idea and fixed how work is partitioned across warps and thread blocks, cutting non-matmul overhead. Later versions target newer hardware features. The idea has not changed since FA1 — the implementations got better at feeding the GPU.

Prefill multiplies a big matrix by a big matrix: many FLOPs per byte moved, so the GPU is well fed. Decode multiplies a matrix by a single vector: the same weights are read to do a fraction of the work. That ratio — FLOPs per byte — is what makes one compute-bound and the other bandwidth-bound.

Batching decode turns vector-matrix products back into matrix-matrix products, amortising the weight read across many requests — the whole reason continuous batching exists. Prefill is already compute-saturated, so batching it adds little and can hurt latency.

A long prefill run in one piece blocks every in-flight decode step. Splitting it into chunks interleaved with decode keeps token streams smooth at a small throughput cost — the same problem prefill–decode disaggregation solves by moving the phases apart entirely.

Memory ≈ 2 · L · S · N_kv · d_head · bytes. The 2 is K and V; L is layers; S is sequence length; N_kv is KV heads. Everything is linear, which is what makes the levers obvious: halve the KV heads (GQA) or halve the bytes (quantise to fp8/int8) and you halve the cache.

Weights are a fixed cost paid once. The KV cache is a per-request cost that grows with every token generated, so as concurrency rises it is the cache — not FLOPs — that fills the card first. This is why serving papers talk about memory almost exclusively.

Four levers, in rough order of how often they are used: fewer KV heads (GQA/MQA), fewer bits per entry (quantisation), better allocation (PagedAttention), and less history to store at all (eviction, compaction, or a sliding window).

Query heads sharing a KV head can still weight it differently — their queries differ — but they can no longer attend to different things. Heads within a group are forced to look at the same set of keys, which removes some of the specialisation multi-head attention exists to provide.

Under TP, heads are split across GPUs. With very few KV heads, either some GPUs hold no KV head at all or the KV must be replicated — so the KV-head count is not a free parameter; it is constrained by how you plan to shard.

Unlike batching or paging, this is architectural. You cannot convert an MHA checkpoint to GQA without retraining or at least a distillation-style conversion step, which is why it is chosen with serving economics already in mind.

Block table ≈ page table; KV block ≈ page; the allocator ≈ the OS memory manager. The difference is that there is no swap — a GPU cannot page to disk cheaply — so when blocks run out the system must preempt a request and recompute its cache later, rather than transparently faulting it back in.

Large blocks mean less bookkeeping but more waste in the final partly-filled block. Small blocks waste less but add indirection and kernel overhead. Typical values are a handful of tokens per block — small enough that internal fragmentation is negligible.

Sharing blocks is what lets parallel samples from one prompt, or beam-search branches, share a prefix instead of duplicating it. When one branch diverges, only the block it writes to is copied.

Decode steps are independent across requests: each one reads its own KV cache and produces one token. Nothing couples request A's step to request B's, so membership can change between iterations without breaking anything. Paged KV is what makes joining and leaving cheap in memory terms.

Adding a request needs KV capacity for its whole future, which you cannot know. Admit too eagerly and you run out mid-generation and must preempt someone; admit too cautiously and the GPU idles. Real schedulers reserve headroom and can preempt-and-recompute the newest request when they guess wrong.

Throughput rises, but per-request latency becomes dependent on load. A p50 measured on an idle server tells you nothing about production. This is why serving benchmarks quote latency at a given throughput, not alone.

A hash map can only answer “have I seen this exact prompt?” A radix tree over token sequences answers the question you actually have: “what is the longest prefix of this prompt that I have already computed?” Branches naturally share their common ancestors, so many conversations sharing one system prompt store it once.

Cached prefixes occupy the same KV memory that running requests need, so the cache needs an eviction policy — typically LRU with a bias toward keeping long, widely-shared prefixes, since those are the expensive ones to rebuild.

Because reuse depends on an exact match, prompt construction becomes a caching problem: stable instructions first, per-request content last, and no timestamps or randomised ordering near the top. It also gives the router a reason to send a request to a specific replica — cache affinity.

A prompt that only ever grows at the tail is a perfect cache workload — every request is a prefix of the next. Context compaction rewrites the middle, which invalidates everything after the edit. In one measured agent run the append-only baseline served 70% of prompt tokens from cache and the compacted run managed 57%. Both techniques are right; they just pull against each other, and raw token counts overstate what compaction actually saves.

The data plane (Envoy sidecars) carries actual traffic and must be fast. The control plane (Istio) never touches a request — it distributes configuration. If the control plane dies, existing proxies keep routing with their last known config; that split is the point of the design.

Mesh defaults assume short, uniform, idempotent requests. LLM calls are long-lived and streamed, so a default timeout can kill a legitimate generation, and a retry can duplicate expensive work. Timeouts, retry policy and streaming support all need revisiting.

Uniform mTLS between services, consistent retry and circuit-breaking, and per-service traffic telemetry — without shipping that logic in every application in every language you use.

Two requests in a queue could be 100 tokens or 100K. Without prompt length, expected output length and KV pressure, a count tells you almost nothing about how long the queue will take to drain.

Nobody knows how long a generation will be, so routers approximate: request type, prompt shape, historical averages per endpoint, or an explicit max_tokens. Getting it wrong is survivable — the router only needs to rank replicas, not predict precisely.

Cache affinity pulls traffic toward one replica; load balancing pushes it apart. Production routers blend them — affinity as a strong preference, overridden once a replica's pressure crosses a threshold.

The KV cache to move is roughly 2 · L · S · N_kv · d_head · bytes — for a long prompt this is gigabytes. On NVLink or InfiniBand that is milliseconds; over ordinary Ethernet it can exceed the prefill you were trying to protect. This is why the technique appears in large deployments and not small ones.

Prefill wants raw compute; decode wants memory bandwidth and capacity. Once separated you can give them different GPU types, different batch policies, even different parallelism strategies — impossible when one pool does both.

The transfer does not have to wait for the whole prefill to finish: each layer's KV can start moving as soon as it is computed, overlapping transfer with the remaining compute.

Splitting a weight matrix by columns lets each GPU compute an independent slice of the output with no communication. Splitting the next matrix by rows means each GPU consumes exactly the slice it already holds, and one AllReduce at the end combines the result. Transformer MLP and attention blocks are deliberately arranged as a column-parallel layer followed by a row-parallel one, so each block costs one collective rather than two.

TP shards weights and activations, so it reduces per-GPU memory the most. PP shards weights by layer but each stage must hold activations for every in-flight micro-batch. EP shards only the expert weights — attention layers are still replicated or TP-sharded.

A large MoE model might run TP=8 inside a node, PP=4 across nodes, and EP across the whole cluster. The dimensions are orthogonal; the art is matching each one's communication pattern to the link that can carry it.

With p stages and m micro-batches, the pipeline spends about (p−1) steps filling and (p−1) draining out of roughly (m + p − 1) total — so the wasted fraction is approximately (p−1)/(m+p−1). Push m well above p and it shrinks; it never reaches zero.

The naive schedule runs all forwards then all backwards, so every micro-batch's activations are held simultaneously — peak memory scales with m. One-forward-one-backward starts backward passes as soon as possible, so activations are freed early and peak memory scales with p instead. It is a memory optimisation, not a speed one.

Giving each GPU several non-contiguous chunks of layers instead of one contiguous block shortens the fill and drain, at the cost of more boundary crossings — worth it when the interconnect can absorb them.

A ring moves each shard around the GPUs in two passes — reduce-scatter then all-gather — so every GPU sends roughly 2(N−1)/N of the data regardless of how many GPUs there are. Bandwidth cost is near-constant in N; only latency grows. That is why AllReduce = ReduceScatter + AllGather is an implementation fact, not just an identity.

Small tensors are dominated by per-hop latency, so tree algorithms win. Large tensors are dominated by bandwidth, so ring algorithms win. NCCL picks per call based on size and topology — which is why a microbenchmark at one message size can badly mispredict real performance.

Communication that cannot be removed can often be hidden: launch the collective on a separate stream and keep computing while it runs. Gradient bucketing in data-parallel training exists precisely so that AllReduce for early layers overlaps with the backward pass of later ones.

Tool schemas, permission policy, sandboxing, retries, timeouts, budget limits, and the stopping rule are all code you write around the model. Improving the model rarely fixes a badly designed loop, which is why agent quality differs so much between products using the same model.

Common conditions: the goal is verifiably met, a step budget is exhausted, no progress across N steps, or a human must approve. Each fails differently — verification is often impossible, budgets truncate mid-task, and no-progress detection needs a definition of progress. Most production agents combine several.

In chat, both sides produce sentences. In an agent, one side produces sentences and the other produces files, logs and pages. Trimming, summarising or referencing observations rather than inlining them is usually the highest-leverage change available.

Some providers attach a full chain of thought to every response. Storing the whole response object in history resends all of it forever. One measured chess run carried 204,027 characters of reasoning against 7,585 characters of actual board observations — games took 31 minutes instead of 3.5. Reasoning is per-turn scratch work, not conversation: keep visible text and tool calls, drop the rest.

A model can emit several calls in one response, all chosen from the same picture of the world. If they only read, that is a free speed-up. If any of them writes, every call after the first was chosen from a world that no longer exists. The chess harness executes the first move and rejects the rest with an explanation, because the opponent has already replied.

Empirically, most failures on repository benchmarks are localization failures, not editing failures: the model could have written the right patch if it had been looking at the right file. This is why retrieval quality and dependency-aware search matter more than raw code-generation ability.

The loop assumes the test result is a signal. A flaky test makes it noise — a passing run may mean nothing. Handling that means running repeatedly, seeding randomness, or isolating the nondeterminism before trusting any result at all.

The agent runs shell commands it wrote itself against a repository it only partly understands. Scoped file access, no network by default, resource limits, and no credentials in the environment turn a destructive mistake into a failed step rather than an incident.

Each instance gives the repository before the fix, the issue text, and hidden tests. The grading is two-sided: FAIL_TO_PASS proves the bug is fixed, and PASS_TO_PASS proves nothing else broke. Only the pair is meaningful — the second is what catches an agent that made things green by deleting a test.

If each step appends d tokens, call t carries about d·t tokens, so the total across T calls is d(1+2+…+T) = d·T(T+1)/2 — quadratic in steps, even though the history itself only grows linearly. Conflating those two quantities is the classic mistake.

Caching removes the cost of re-reading history that is still present; the window fills at exactly the same rate. Compaction removes the history itself. If you are hitting a cost ceiling, cache. If you are hitting the context limit, only compaction helps.

Compaction discards silently, and you find out several steps later. Practical mitigations: keep a structured summary the agent maintains deliberately rather than a free-text one, keep artifact references rather than artifact contents, and keep the failed approaches — re-trying a ruled-out fix is the most common drift symptom.

In one measured coding run the agent spent 463,000 prompt tokens against 13,000 completion tokens — roughly 35 to 1. The bill is for re-reading, not for thinking, which is why every lever in this section targets the prompt rather than the generation.

Pure replay rebuilds state by re-running an event log — exact, but slow and dangerous if events had side effects. Pure snapshots are fast but lose everything since the last one. Production systems take periodic snapshots and replay only the log since, which bounds both recovery time and loss.

The durable record has to exist before the side effect, not after — otherwise the crash window between acting and recording is exactly the case you cannot recover from. Write “about to charge, key=abc”, then charge, then reconcile.

In: goal, plan, step status, pending action and approval state, artifact references, environment handles. Out: the full transcript, and anything reconstructible. A checkpoint that grows with the conversation has become another context problem.

A run can reach a correct final state while a quarter of its steps produced no action at all. Grading only the outcome misses that entirely. Inspect trajectories for the behaviour you care about — steps with no tool call, repeated identical actions, retries after a permanent failure — and A/B the harness, not just the model, because the interface is as much of the system as the weights are.

A tool is an action with side effects the model may choose to call. A resource is addressable data the host can read into context without the model deciding anything. A prompt is a reusable template the server offers. Conflating tools and resources is the most common misreading of the protocol.

A host can ask a server what it offers at runtime and present those capabilities to the model, instead of having them hard-coded at build time. That is the difference between a standard and a convention — and it is what turns N×M integrations into N+M.

Every connected server widens what the agent can reach, and tool descriptions are text the model reads — so a malicious server can attempt to influence behaviour through them. Permission checks, approval gates and treating server-supplied text as untrusted data remain the host's responsibility.

Documents already exist — nobody decided they were worth writing for the agent's benefit. Memories must be chosen, which adds a whole problem RAG never has: what, when, and at what granularity to save, and when to overwrite something now false.

The model can only use what enters the context. Chunk boundaries that split a function in half, or an embedding that matches on topic rather than intent, fail silently — the answer looks confident and is simply missing the relevant material. Hybrid retrieval (lexical + semantic) exists because each covers the other's blind spot.

A wrong conclusion written once gets retrieved forever, and each retrieval makes it look better established. Practical defences: store what was observed rather than what was concluded, timestamp entries, and allow later evidence to supersede rather than accumulate.

Retrieval assumes the query resembles what you need. Much of what an agent must not forget fails that test: “I already tried X and it failed” has to be present regardless of what the next step is about, and nothing in that step’s text would retrieve it. Summaries also preserve order and causality, which retrieved fragments lose — and retrieval hands back the raw output you were trying to shed. Retrieval does belong in coding agents, but on the codebase, not the transcript.

Once the model has emitted a tool call, the harness has no way to un-emit it. Every path out of the dispatcher — success, bad JSON, unknown tool, wrong argument type, non-zero exit — has to produce exactly one observation carrying that id. Treating it as a total function rather than one that can throw is the design.

Step t resends everything from steps 1..t−1. If each step adds d tokens, the sum across T steps is d·T(T+1)/2. Completion tokens grow linearly; prompt tokens grow as the square. That is why one measured run showed 463k prompt against 13k completion — a ratio of about 35 to 1.

Some models attach hidden reasoning to every response. Storing the whole response object in history resends all of it forever. In one run that was 204,027 characters of reasoning against 7,585 characters of actual observations — the prompt was 96% scratch work. Keep visible text and tool calls; drop the rest.

Each compaction is handed the previous working memory along with the new prefix, and told to drop what later steps made obsolete. A memory that only appends recreates the exact problem it was built to solve — it just grows more slowly.

The summariser reads a readable transcript, not raw JSON message objects. It is a one-off rendering cost that buys a markedly better summary, because the model is being asked to read a conversation rather than parse a data structure.

Two of five summaries in the measured run came back with finish_reason length — they hit the 1,200-token output cap mid-sentence and the tail was lost. Nothing failed, nothing logged an error, and the run still succeeded. Either raise the cap or keep more recent steps verbatim, but first make the truncation visible.

On a transport failure, ask the environment whether it is still alive and abort with a clear message if it is not. Add a five-consecutive-failure backstop for the case the probe itself succeeds over a dead tunnel. Illegal moves never increment the counter, and any success resets it.

Too eager to recover gives an infinite loop billed by the token. Too eager to crash gives brittleness — a single flaky network call kills a run that would have succeeded on retry. Both are error-policy bugs; only the first one is expensive enough to notice immediately.

A deliberate abort is only cheap if the run can resume. Terminating on a dead sandbox and then starting the whole task over is not much better than retrying. The termination condition and the checkpoint are two halves of the same design.

Truncating a 120,000-character result to its first 10,000 throws away the end of the stack trace — usually the line that names the failure. Keeping the first and last 4,900 with an elision notice between them costs the same tokens and keeps the part that identifies the problem.

Progressive disclosure is usually discussed for skills or documents, but a tool schema with a long description is the same cost in the same place. Anything carried in every request to describe a capability the agent might not use is a candidate for the split.

If a single step blew up the prompt, you need truncation. If the prompt crept up over forty steps, you need compaction. If the prompt was already large before the agent did anything, you need progressive disclosure. The symptom tells you which.

The saving in round trips is the obvious part. The less obvious part is that procedures too long to spell out as individual calls become possible at all — a deeper search, a retry loop, a multi-stage transformation. The unit of work stops being the thing that limits the algorithm.

A snippet may have changed live state several times before returning. The harness cannot assume its cached picture is still accurate, so after every snippet it re-reads the environment. Skipping this is exactly the stale-state bug, just at procedure scale.

A sequence of individual tool calls leaves a readable trail: you can see which call failed and what it returned. A failed snippet gives you one opaque observation. That is a genuine argument for keeping ordinary tool calls as the default and reaching for code only when the work is loop-shaped.

No, and the confusion matters. A query here is an attention query vector for one token: the model projects a token's hidden state into a query and uses it to decide which past information should flow in. So T = 8 means the kernel is processing eight query tokens — not eight user requests. That is why the benchmark's tiny T is a hardware problem rather than a product decision.

No. The kernel assumes the sparse indices already exist; selection happens upstream in the indexer. Saying this clearly scopes the work: everything we did is about performing attention over an arbitrary scattered index list efficiently, not about deciding which tokens are relevant.

Only on chip, and only for the current tile. The source stays scattered in HBM forever — gather pulls the currently needed subset into a small shared-memory tile where the layout is regular enough for Tensor Cores, and then pays the scatter again for the next tile. Which is why hiding the gather (cp.async) and having more warps to hide it behind (Split-K) are both worth doing.

QKᵀ gives attention scores; softmax(QKᵀ) gives attention weights; softmax(QKᵀ)V is the attention output. The mental model worth carrying: QK decides where to look, PV collects the information.

The threads in a block need a single shared-memory allocation and a single __syncthreads() barrier domain. Both are SM-local resources. Splitting a block across two SMs would mean shared memory that spans two physical memories and a barrier across two schedulers — which is precisely the cost the block abstraction exists to avoid. The indivisibility is the price of shared memory being fast.

Each SM has its own shared-memory resource (up to 228 KB on B200), and CUDA carves a private region of it for every resident block. Ordinarily blocks never see each other's (clusters, below, are the one exception). The practical consequence shows up everywhere in this project: more shared memory per block means fewer blocks resident per SM, which means fewer warps available to hide a stall.

A warp that issues a scattered load and misses cache stalls for hundreds of cycles. If another warp is resident and ready, the scheduler issues from that one instead and the execution units stay busy. With regular, coalesced access this matters less; with a gather from arbitrary HBM addresses it matters a lot — which is why the final split kernel doubled its warp count.

Since Hopper, blocks can be grouped into a cluster that the hardware guarantees to schedule together on neighbouring SMs. Blocks in a cluster can read and write each other's shared memory directly — distributed shared memory — instead of round-tripping through global memory. On B200 the portable cluster size is 8 blocks, with 16 available as a non-portable option. It refines the picture rather than overturning it: each block still lives on exactly one SM; the cluster just lets neighbours reach across.

When a thread needs more registers than it is allowed (at most 255), or indexes a per-thread array dynamically, the compiler spills to “local memory” — which is private to the thread but physically lives in global memory, cached in L1/L2. A spill-heavy kernel quietly turns its fastest storage into its slowest. nvcc -Xptxas -v reports spill bytes; it is worth checking whenever a kernel is mysteriously slow.

It is tempting to treat high occupancy as the goal. It is not — it is one way to buy latency hiding. A kernel that is genuinely compute-bound gains nothing from more resident warps, and a kernel that spends shared memory on deeper buffering may be better off with fewer. The question is always which resource is actually idle.

Fewer bits per number means less data movement and much higher Tensor Core throughput, but accumulating thousands of products in low precision loses bits fast. The standard shape — low-precision multiply, higher-precision accumulate — keeps the throughput while holding the running sum stable. It is the same reason FlashAttention accumulates its softmax state in FP32.

The competition required matching the BF16 reference within abs_err under 0.1. That is a hard wall: it rules out FP8 shortcuts and approximate softmaxes. Worth mentioning unprompted — it turns “we used BF16” from a buzzword into a decision made inside a constraint.

They solve different bottlenecks and get conflated constantly. WMMA maps matrix math onto Tensor Cores. FlashAttention-style tiling stops intermediate score and probability matrices from being written to HBM at all. You can have either without the other; this kernel has both, for different reasons.

It is easy to call v1 wasted effort. It was not — until the arithmetic was fast, you could not see that memory was the limit, and the measurement that showed 600 µs after a large compute speedup is exactly the evidence that pointed at the gather. Optimisation is mostly a sequence of measurements that relocate the bottleneck.

In steady state the stage cost is max(T_load, T_compute), not the average. If the gather costs 10 and the compute costs 2, you hide 2 and the other 8 are still on the critical path. The measured 600 → 470 µs says the gather was much larger than the compute — which is the diagnostic that pointed at the next optimisation.

Deeper pipelining is possible, but each stage costs shared memory per block, and shared memory caps how many blocks stay resident per SM. Past a point you are spending occupancy — which is itself a latency-hiding mechanism — to buy more overlap. The two levers compete for the same budget.

An ordinary global load occupies the issuing warp until it returns. cp.async issues a global→shared copy that the warp does not wait on, so the warp can continue issuing MMA work and synchronise on the copy later. Without the async form, two buffers would buy you nothing — the warp would still be blocked.

If the kernel is saturating memory bandwidth, overlap does not help — there is no idle transfer capacity to fill, and the answer is to move fewer bytes. Here the problem was latency on scattered accesses rather than raw bandwidth, which is why hiding it worked at all.

Split-K is often explained as “the working set doesn't fit in shared memory, so split it”. That is not what is happening here. A single block could already stream all 2,048 entries with tiling and double buffering — nothing overflowed. The purpose is to expose more independent block-level parallelism. Getting this right is the difference between reciting the name and understanding the decision.

Split-K creates S partial states per query, writes them to global scratch, and requires a second kernel to read them back and merge. Total work goes up. Wall-clock latency goes down, because far more of the machine is computing simultaneously. Latency and throughput are different quantities and this is the cleanest example of them diverging.

More splits means more parallelism and more reduction overhead. At small S the GPU is underfilled; at large S the reduce kernel dominates and each block's tile is too small to amortise its own setup. Somewhere in between is the minimum, and where it sits depends on the workload — hence v4.

“The work does not divide into enough independent pieces to fill the hardware” is the same complaint as a pipeline bubble or an unbalanced tensor-parallel shard. The fix is always to re-cut the problem along a different axis and pay a combination cost. Here the axis is the KV index list and the cost is an online-softmax reduction.

mᵢ is the chunk's maximum score, lᵢ = Σ exp(s − mᵢ) is its un-normalised denominator, and oᵢ = Σ exp(s − mᵢ)V its un-normalised output. Rescaling by exp(mᵢ − m) onto a common maximum m converts any chunk's statistics into the global frame, so the merge is exact and does not depend on the order chunks arrive in.

It is what keeps exp from overflowing. Attention scores can be large; exp(score) in FP32 overflows well before the scores become unusual. Carrying a running max and exponentiating only differences is the numerically stable form — the same reason single-kernel softmax implementations do a max pass first.

FlashAttention carries (m, l, o) from tile to tile inside one block so the S×S score matrix never reaches HBM. Split-K carries the identical state from block to block so 32 processors can reconstruct one softmax. Recognising that these are the same trick at different scales is a good thing to say out loud.

The correctness bar was matching the BF16 reference within abs_err 0.1. An approximate merge — averaging, or dropping small chunks — would have failed outright. Split-K is only available as an optimisation because an exact merge exists.

The reduce kernel ran 4 warps over 16 heads — four serial passes. v4 uses 16 warps (512 threads), one per head, in a single pass. The split kernel went from 4 warps (128 threads) to 8 (256), so more independent warps are resident to cover a stalled scattered read. Neither is clever; both are about not leaving units idle.

A sweep tells you which configuration is fastest, not why. That is fine for shipping and unsatisfying for understanding — the honest answer in an interview is that the measurement decides and the profiler explains. If you can say which resource the winning configuration stopped wasting, you are ahead of the sweep.

It costs compilation and benchmarking time, which is acceptable offline and awkward when shapes arrive at runtime. The usual production compromise is a cached table keyed by shape, with a heuristic as the fallback for anything unseen — which is roughly what v3 was, used correctly.

It is the same lesson as least-outstanding-requests beating round-robin in a load balancer: a model of the system is a simplification, and where the simplification is wrong the measurement is not. Reaching for a formula is the right instinct; refusing to check it against the hardware is not.

A single neuron's decision boundary is the set where w·x + b = 0 — a line — and XOR's two classes sit on opposite corners of a square, so no line puts both +1 corners on one side without one −1 corner. Two hidden ReLUs, h₁ = ReLU(x₁ + x₂) and h₂ = ReLU(x₁ + x₂ − 1), map the four corners to (0,0), (1,0), (1,0), (2,1), which are separable.

The rule w ← w + t·x is gradient descent on the loss max(0, −t·(w·x + b)) with learning rate 1: the gradient of that loss on a mistake is exactly −t·x. Swap the step output for a sigmoid and the hinge for cross-entropy and you have logistic regression — same neuron, differentiable loss, and the gradient now carries how wrong, not just whether.

“Why do we need nonlinear activations?” Because a stack of linear layers is one linear layer: W₂(W₁x + b₁) + b₂ = (W₂W₁)x + b′. The activation between them is the only thing that lets depth draw anything but a line.

The TensorFlow Playground lets you watch a tiny network bend the plane on exactly these datasets, XOR included, live in the browser. Rosenblatt’s 1958 paper and Minsky & Papert’s Perceptrons (1969) are the historical bookends; the Wikipedia article has the convergence theorem and the original learning rule.

Sigmoid and tanh saturate: for large |z| their slope is nearly zero, so gradients vanish through deep stacks. ReLU's slope is exactly 1 wherever it is on, so signals pass through unchanged, and it costs one comparison. Its flaw is units that are off for every input (dead ReLUs); GELU and SiLU, used in transformers, smooth the corner so there is always some slope.

With a sigmoid output and squared error, the gradient at the logit is (ŷ − y)·ŷ(1 − ŷ), which goes to zero exactly when the network is confidently wrong (ŷ near 0 or 1) — the worst time to stop learning. With cross-entropy the σ′ factor cancels and the gradient is just ŷ − y: largest when most wrong.

3Blue1Brown’s neural-network series (chapter 1) animates exactly this forward pass on handwritten digits. Stanford’s CS231n notes draw the layer-as-matrix picture used here.

For a batch of B examples, X is B × d_in and the layer computes X W + b with W of size d_in × d_out — one matmul of B·d_in·d_out multiply-adds. A transformer's MLP has d_in = 4096, d_out = 16384; with B = 4096 tokens that is 2.7 × 10¹¹ multiply-adds in one line. The forward pass is the reason the infrastructure notes exist.

ReLU won in 2012 because its slope is 1 and it costs a comparison; the price is dead units and a kink at zero. GELU (z·Φ(z)) and SiLU (z·σ(z)) are smooth, never have slope exactly zero, and in practice give slightly better transformers at the same cost — GPT-2/3 and BERT use GELU, Llama uses SiLU inside a gated unit (SwiGLU).

Hidden activations are chosen for their slope. The output activation is chosen to match the loss: sigmoid + binary cross-entropy for one probability, softmax + cross-entropy for a distribution, identity + squared error for a regression. In each matched pair the gradient at the logits collapses to prediction − target — the same ‘gradient = error’ every time.

The table on Wikipedia’s activation-function page plots every function with its derivative. Stanford’s CS231n notes (neural networks part 1) walk through sigmoid, tanh, ReLU and leaky ReLU with the same ‘what does it do to the gradient’ framing used here. The GELU paper is Hendrycks & Gimpel, arXiv:1606.08415.

∂L/∂ŷ = −1/ŷ and ∂ŷ/∂u = ŷ(1 − ŷ); multiply and the ŷ cancels, leaving ŷ − y. The same cancellation happens for softmax + cross-entropy in a classifier with many classes: the gradient at the logits is simply (probabilities − one-hot). That is why those two are always paired, and why the error signal at the output of a classifier is literally the error.

For a layer h = W a, backward needs two products: ∂L/∂W = (∂L/∂h) aᵀ to update the weights and ∂L/∂a = Wᵀ (∂L/∂h) to keep going. Each is the same size as the forward matmul, so backward ≈ 2× forward, and a training step ≈ 3× a forward pass — the figure behind every FLOPs estimate for training.

3Blue1Brown’s Backpropagation calculus draws the chain-rule tree for a one-neuron-per-layer network, the same shape as the chain panel here. Karpathy’s micrograd (and its two-hour video) builds backprop in 100 lines of Python on a graph just like the one on the autodiff page.

Each layer multiplies the signal by its weight matrix and by the activation's slope. Across 50 layers that product either shrinks toward zero (sigmoid, small weights) or grows without bound. Residual connections add an identity path so the product always has a 1 in it; LayerNorm keeps the scale of activations fixed; careful initialisation starts the products near 1.

Both use the same local derivatives. Forward mode pushes a perturbation of one input through the graph (a Jacobian–vector product) and gets the sensitivity of every output to that one input; reverse mode pulls a weight on one output back through the graph (a vector–Jacobian product) and gets the sensitivity of that one output to every input. Training has one output (the loss) and billions of inputs, so reverse mode wins by a factor of billions.

Weights, their gradients, the optimizer state (two more copies for Adam), and the saved activations — which for a transformer are every layer's input, attention scores and MLP intermediate for every token in the batch. The activations usually dominate at long sequence lengths. Activation checkpointing keeps only the layer boundaries and recomputes the rest during backward, trading ~30% more compute for a large memory cut.

PyTorch’s Autograd mechanics note is the authoritative description of what gets saved and when; Baydin et al., Automatic differentiation in machine learning: a survey, has the canonical forward-vs-reverse mode tables. Karpathy’s micrograd is the smallest working implementation worth reading end to end.

PyTorch records the graph fresh on every forward call, so Python control flow (if, loops, early exit) just works, at the cost of re-recording each time. Static-graph systems (TensorFlow 1, JAX's jit, torch.compile) capture the graph once and can optimise it — fuse kernels, plan memory — but need the structure to be fixed. Modern practice is eager for development and compiled for the training run.

Doubling B halves the gradient variance, so the same per-epoch progress survives doubling η too — the linear scaling rule (Goyal et al., arXiv:1706.02677, trained ImageNet in an hour on 8k-image batches with it, plus a warm-up). It stops working past the critical batch size, where the gradient is already accurate and extra examples per step buy nothing; beyond that only more steps help.

Keskar et al. (arXiv:1609.04836) observed that large-batch training lands in sharper minima and generalises worse, and attributed it to the missing noise. The picture is widely used and well supported but still argued about — later work showed the gap can mostly be closed by scaling η and training longer. In an interview, say ‘tends to’, not ‘always’.

The loss-landscape renders everyone has seen — the spiky ResNet-56 surface that skip connections smooth out — are from Li et al., Visualizing the Loss Landscape of Neural Nets. The optimizer-race animations (SGD, momentum, NAG, Adagrad, Adadelta, RMSprop racing through a saddle and a valley) are Alec Radford’s, reproduced in CS231n notes part 3, which also covers batch size, learning-rate schedules and gradient checks. Sebastian Ruder’s overview of gradient descent optimization algorithms is the standard written tour.