The Kernel Level: Attention and MoE

Picture yourself profiling a 70B model on an H100 for the first time. You open the trace, and something is off: attention at long context takes seconds, yet your FLOP-counting from Lecture 1 says the math is milliseconds. Then you look at a decode step and find hundreds of tiny kernels separated by... nothing. Gaps. Idle silicon. Two leaks stand between your paper derivation and the metal: bytes that should never have left the chip, and time lost between kernels. This lecture is the hunt for both.

Lectures 1–6 decided what to compute and where; today is how one GPU actually computes it: which bytes live at which level of the memory hierarchy at which instant, and how hundreds of kernel launches are stitched into one decode step. Two kernel families carry almost all the load. The attention main line is FlashAttention → FlashInfer → FlashMLA → the sparse DSA chain; the MoE main line is grouped GEMM → DeepGEMM → DeepEP. Running through both is an execution-style evolution: manual launches → CUDA graphs → Mega MoE's communication–compute fusion. Every step is a response to a derivable number; deriving those numbers is our business today.

The Design Space: Four Levels of Memory

An H100 streaming multiprocessor (SM) presents you with a four-level hierarchy. Only two levels are under your exact control: registers and shared memory (SMEM). That is why kernel design is residency policy: which tiles are pinned in SMEM, which are streamed, and which are never materialized at all. Every byte crossing HBM costs \(1/\beta\) seconds; bytes that stay on chip are, to first order, free.

Three Hopper mechanisms changed what a good residency policy even looks like: TMA (async block-wise global↔shared copies; one thread issues, the hardware does the address arithmetic, think block-table lookup for paged KV, and signals an SMEM mbarrier); WGMMA (matrix multiply issued asynchronously per warpgroup, operands straight from SMEM); and thread-block clusters with distributed shared memory (blocks read each other's SMEM: a tile-exchange pool larger than 228 KB without touching L2/HBM). With copies and MMAs both asynchronous, the natural post-Hopper kernel is no longer "all warps do the same thing": producer warps run TMA for tile \(t+1\), consumer warpgroups run WGMMA on tile \(t\), epilogue warps drain tile \(t-1\), all handing off through mbarriers. This producer/consumer/epilogue skeleton is shared by FlashAttention-3, DeepGEMM, and Mega MoE. It keeps the tensor cores fed across the ~500-cycle HBM latency by decoupling "move bytes" from "multiply bytes". Two cousins to remember: Blackwell (SM100) adds tcgen05 with TMEM and native FP4 datapaths (they return below), and persistent kernels fix the grid at the SM count and pull tiles from a task queue: no re-launch cost, no wave quantization, and the freedom to interleave tiles of different subproblems onto free SMs.

Naive Attention: the \(O(L^2)\) You Should Never Write

Here is the easiest attention kernel you can write. If you have used eager PyTorch, you have written it: compute the full score matrix \(S = QK^\top\), softmax it into \(P\), multiply by \(V\). You can probably guess this is not the best way. But let us count exactly why not, because the number is spectacular. Assumptions: prefill of a length-\(L\) sequence, one layer, \(H_q = d_{model}/d_h\) query heads, \(n_{kv}\) KV heads, BF16 (\(b_{dtype} = 2\) B). We count per query head: GQA shares \(K, V\), but every query head owns its own \(S\) and \(P\).

Summing over heads, traffic and arithmetic intensity of the naive implementation are \[ T_{naive} = H_q\left(4L^2 b_{dtype} + 4Ld_h b_{dtype}\right) \approx \frac{4L^2 d_{model}\, b_{dtype}}{d_h}, \qquad I_{naive} \approx \frac{d_h}{b_{dtype}} = \frac{128}{2} = 64\ \text{FLOP/byte}. \] (The FLOPs, for contrast, are \(\approx 4L^2 d_{model}\) per layer: Lecture 1's attention term.) And 64 sits far left of the H100 ridge \(\pi/\beta \approx 295\): attention executed naively is bandwidth-bound even at prefill sizes. The \(O(L^2)\) traffic, not the \(O(L^2)\) FLOPs, is the binding constraint. This is the single most important sentence about attention kernels.

But Softmax Wants the Whole Row...

Fine. Never write \(S, P\) to HBM. Just one small obstacle, and you may have spotted it: softmax needs the whole row before any output is final, because \(\mathrm{softmax}(s)_j = e^{s_j} / \sum_t e^{s_t}\) couples all KV positions through the denominator. Tiles are small; the row is long. How do you normalize a sum you haven't finished reading? Think first (is a running denominator even possible?), then read on.

The trick is online softmax: per query row, keep a running max \(m\) and running normalization \(\ell\), and correct past partial results by exponential rescaling whenever \(m\) grows. Process the score row \(s = qK^\top \in \mathbb{R}^L\) in blocks \(s^{(1)}, \dots, s^{(T)}\) of length \(B_c\) (with \(V\) blocked conformably); from \(m_0 = -\infty,\ \ell_0 = 0,\ O_0 = 0\): \[ m_i = \max\big(m_{i-1},\ \mathrm{rowmax}(s^{(i)})\big), \qquad \ell_i = \ell_{i-1}\, e^{m_{i-1} - m_i} + \sum_j e^{s^{(i)}_j - m_i}, \qquad O_i = O_{i-1}\, e^{m_{i-1} - m_i} + e^{s^{(i)} - m_i} V^{(i)}, \] with final output \(O_T / \ell_T\). This unnormalized form defers all division to the end (a normalized variant divides each step; the two are identical up to carrying \(1/\ell_i\) inside \(O\)).

Tiling turns the recurrence into a traffic theorem: per KV block the kernel touches \(2 B_c d_h\) bytes of HBM (the \(K_j, V_j\) tile) and zero bytes for \(S_j, P_j\), which exist only in registers/SMEM. Summed: \(Q, K, V\) read once, \(O\) written once. Exactly the \(O(Ld)\) traffic we counted.

On Hopper, FA3 runs this loop on the warp-specialized pipeline: producer TMA warps stream the next \(K_j, V_j\) while consumers run WGMMA, and element-wise softmax overlaps the async MMA (the "GEMM–softmax overlap" pushing FA3 toward the compute ceiling; FA3 also adds FP8). The family history in one line: FA1 established online softmax + tiling; FA2 re-parallelized along the Q sequence (each block owns a Q tile and scans all KV, filling all 132 SMs even at long \(L\) and small batch); FA3 rebuilt the kernel for Hopper. Same mathematics every time. What evolves is the residency and overlap policy.

Decode Starves the SMs: Split-KV, the LSE Merge, and One Latent for 128 Heads

Now decode arrives and breaks the map. At decode \(q_{len} = 1\): exactly one query row per request, so FA2's "split along Q" yields only batch-size-many blocks, often far fewer than 132 SMs. You can probably guess where the parallelism must come from: the only long axis left is the KV length itself.

Split-KV (a.k.a. flash-decoding) divides the \(L\) KV positions into \(S\) segments; each segment runs the ordinary online-softmax recurrence in its own block and produces a partial triple \[ m_s = \max_{j \in s} s_j, \qquad \ell_s = \sum_{j \in s} e^{s_j - m_s}, \qquad O_s = \sum_{j \in s} e^{s_j - m_s} v_j. \] A small reduction kernel then merges the \(S\) partials. Let \(m^* = \max_s m_s\); every partial is deflated to the common scale exactly as in the online recurrence: \[ \ell^* = \sum_s \ell_s\, e^{m_s - m^*}, \qquad O^* = \frac{\sum_s \ell_s e^{m_s - m^*}\, O_s^{norm}}{\ell^*}, \qquad O_s^{norm} = O_s/\ell_s. \] Substitution collapses the merge into the global sum: the segments partition \([L]\), so this equals full-KV softmax by the same cancellation. Two equivalent forms of one correction: define each segment's log-sum-exp \(\mathrm{LSE}_s = m_s + \log \ell_s\); the merge weight of segment \(s\) becomes \(e^{\mathrm{LSE}_s - \mathrm{LSE}^*}\), a softmax over segment LSEs applied to partial outputs (verified numerically to \(6\times 10^{-16}\)).

A kernel alone is not a serving system. FlashAttention assumes dense contiguous Q/K/V and one shape per call; an engine needs FlashInfer's contract: native paged-KV support (the kernel indexes KV through Lecture 2's block table, no contiguous copy needed), ragged batches carried as metadata, a unified prefill/decode/append interface, JIT-compiled variants (head dim, RoPE, soft-capping, sliding window), and plan/run separation, where a host-side plan inspects the batch's length distribution and decides tile partitioning and load balance, keeping variable-length batches schedulable without device-side sync. Both vLLM and SGLang use FlashInfer as an attention backend: it is where Lecture 2's paging meets this lecture's tiling.

And then there is MLA, whose decode shape breaks standard kernel assumptions. Recall Lecture 2: in absorbed mode, per-head projections fold into \(Q\), so attention runs directly against the compressed cache: each token contributes one 576-element latent (512-dim latent + 64-dim RoPE key) shared by all 128 query heads. The decode GEMM is \[ \underbrace{O}_{128 \times 576} = \mathrm{softmax}\Big( \underbrace{Q_{lat}}_{128 \times 576}\, \underbrace{K_{lat}^{\top}}_{576 \times L}\Big)\, \underbrace{V_{lat}}_{L \times 576}, \] attention with 128 heads packed as GEMM rows against a single K-dim-576 cache. The naive treatment (128 independent heads) re-reads the same latent cache 128 times: 72 KB per token instead of 576 B. That forfeits exactly the bandwidth MLA was designed to buy. FlashMLA loads each block of latent KV from HBM once and lets all 128 packed heads consume it from on-chip memory.

Last stop on the attention line: sparsity. FlashMLA additionally accepts an indices tensor (batch, qlen, topk) and gathers the selected tokens from paged KV directly into SMEM by index, materializing nothing intermediate. This is the execution endpoint of DSA (Lecture 3). The full per-step chain in DeepSeek V3.2/V4 is three kernels:

MoE: 256 Experts, 8 Tokens Each

Enough attention. The other half of the decode step is the MoE FFN. Under expert parallelism (Lecture 6) each GPU owns \(E_{local}\) experts and, every step, must run one GEMM per local expert. Two facts conspire against you. First, the per-expert token counts \(M_i\) are unequal (routing is data-dependent) and known only at runtime. Second, and this is the killer: at decode they are tiny. With \(B = 256\) requests and top-8 routing over 256 experts, the mean is \(M_i = 8\).

Why not one launch per expert? Two failures. Launch overhead multiplies by the expert count: \(E_{local} \times\) (Linear1 + SwiGLU + Linear2) at a few microseconds each (we will price a launch very soon). More fundamentally, tile quantization: tensor-core GEMMs compute in tiles of, say, 128 rows, so a GEMM with \(M_i = 8\) useful rows still pays for full tiles: \[ \eta_{tile} = \frac{\sum_i M_i}{\sum_i \lceil M_i / T_M \rceil\, T_M} = \frac{\text{useful FLOPs}}{\text{executed tile FLOPs}}, \qquad M_i = 8,\ T_M = 128 \Rightarrow \eta_{tile} = 6.25\%. \] The tensor cores execute 16× more work than is useful (random routing barely helps: a multinomial draw of 2048 assignments over 256 experts still gives \(\eta \approx 6.25\%\)). At prefill the same shape is fine: 16K tokens × top-8 over 256 experts gives \(M_i \approx 512 = 4 \times 128\), so \(\eta \approx 100\%\). The MoE GEMM efficiency problem is a decode-phase, small-\(M\) problem. That is why its solutions (persistent kernels, masked layouts, fusion) all target the decode path.

Grouped GEMM is the fix: one kernel schedules all subproblems \(\{(M_i, N, K)\}_{i=1}^{E_{local}}\). A routine builds a tile-level task list (every (expert, row-tile, col-tile) triple) and a persistent kernel pulls tiles from the queue. Tiles of different experts interleave freely across SMs, so small-\(M_i\) GEMMs no longer strand SMs at per-expert wave boundaries: the SM pool sees one long queue of same-shaped tile tasks. Tile quantization within each expert remains (pooling does not shrink tiles), but wave quantization and launch multiplicity are gone. One further design choice is a phase decision:

• Contiguous layout (prefill). Tokens sorted by expert and packed tight; each subproblem gets its exact \(M\). The sort costs one pass, amortized over prefill's many tokens; the GEMM sees dynamic but tight shapes, no wasted rows.
• Masked layout (decode). Each expert gets a fixed-capacity slot of \(C\) rows; the real count rides in a count array, unfilled rows computed-then-discarded. Shapes are static every step, precisely what CUDA-graph capture requires. The price is padding: \(C = 32\) with realized \(M_i = 16\) gives \(\eta_{tile} = 50\%\). DeepGEMM provides this masked variant and accepts an expected-\(M\) hint to size its persistent grid.

The library executing this is DeepGEMM: FP8/FP4/BF16 GEMMs, grouped GEMM in both layouts, the DSA indexer, and the Mega MoE kernels, all JIT-compiled. Two of its choices deserve a close look. Fine-grained FP8 scaling: naive FP8 keeps one scale per tensor, and with ~4 mantissa bits (E4M3) the dynamic range is poor; DeepGEMM gives every \(1 \times 128\) activation strip and every \(128 \times 128\) weight block its own scale, applied at dequantization, plus two-level accumulation: FP8 tensor-core accumulators are not precise over long \(K\) reductions, so partial sums are periodically promoted to higher precision on CUDA cores. Bolted-on FP8 with per-tensor scales never reaches usable accuracy: "just cast to FP8" is not a strategy for V3-class models. JIT as philosophy: no precompiled template zoo. Each concrete \((M, N, K)\), layout, and precision gets a specialized, cached kernel with tiles, unrolls, and pipeline depth matched to it. The cost is first-call compile latency: profile a JIT kernel cold and you measure the compiler, not the kernel. Published: ≈ 1550 TFLOPS FP8 on H800 (≈ 78% of the ≈ 1979 TFLOPS FP8 peak, a high MFU for a fine-grained-scaled GEMM), built on TMA + warp specialization on Hopper, tcgen05/TMEM + native FP4 on Blackwell.

Completing the MoE main line, DeepEP is the communication library behind Lecture 6's all-to-alls, in two phase-specific modes. Normal mode (prefill) aggregates tokens intra-node over NVLink (≈ 900 GB/s), then forwards cross-node over RDMA (≈ 50 GB/s per GPU), minimizing expensive inter-node bytes. Low-latency mode (decode) skips the aggregation hop and sends pure RDMA point-to-point, on the order of a hundred microseconds, with a hook mechanism: dispatch/combine return right after initiation, reception runs in the background, and overlap interleaves at micro-batch granularity (while micro-batch 1's tokens are in flight, micro-batch 2's experts compute). Everything is GPU-initiated RDMA (the IBGDA/NVSHMEM line: GPU threads ring the doorbells, no CPU proxy), and precision follows Lecture 6's contract: dispatch FP8, combine BF16.

Between the Kernels: Launch Overhead, CUDA Graphs, and Fusion

So far we optimized inside kernels. Zoom out one level: a forward pass of a 60–80-layer model runs a dozen-plus kernels per layer, hundreds per step. In prefill, kernels run hundreds of microseconds and launch cost is noise; in decode, kernels run microseconds, and now the stuff between kernels is the leak. Let a decode step execute \(n_k\) kernels with total GPU execution time \(T_{exec}\), each launch costing \(t_{launch}\) of CPU submission/queueing plus \(t_{gap}\) of inter-kernel GPU idle. If launches don't fully overlap execution (small batches: the GPU drains faster than the CPU enqueues), \[ T_{step} = T_{exec} + n_k\,(t_{launch} + t_{gap}). \] Substitute a 70B-class model: \(80 \times \approx 12\) kernels per layer gives \(n_k \approx 960 \approx 10^3\), and with \(t_{launch} + t_{gap} \approx 2\text{–}5\,\mu s\): \[ n_k (t_{launch} + t_{gap}) \approx 1000 \times (2\text{–}5\,\mu s) = 2\text{–}5\ \text{ms}, \] against a small-batch decode step of ≈ 10 ms. 20–50% of the step is launch machinery, not math. No kernel-level optimization can reclaim it; the overhead lives between kernels.

CUDA graphs capture the whole step's kernel sequence (parameters, dependencies, stream structure) into a replayable object; each subsequent step is one graph launch, and both overhead terms die together. The catch is static shapes and addresses: inside a captured graph every buffer pointer and shape is frozen. Three engine mechanisms jointly satisfy the freeze:

1. Batch-size buckets. Capture several graphs for different batch sizes, pad to the nearest bucket at runtime. With bucket width \(w\) and batch size uniform on \([1, B_{max}]\), mean padding is \((w-1)/2\) tokens per step: \[ \mathbb{E}[\text{pad}]/\mathbb{E}[B] \approx \frac{w}{2\,\mathbb{E}[B]}, \qquad w = B_{max}/K\ \text{for}\ K\ \text{buckets.} \] At \(B_{max} = 256\): \(K = 4\) buckets (\(w = 64\)) waste 24.5% of compute on average; \(K = 16\) wastes 5.8%; \(K = 32\) wastes 2.7%. Graph count buys back padding at a \(1/K\) rate, at the price of capturing and storing \(K\) graphs.
2. Paged indirection. The KV cache grows every step: an inherently dynamic address. The graph freezes not the KV pages but the block-table pointer; the pager updates the table each step and kernels dereference pages through it at fixed addresses. Paging is not only memory management; it is what makes growing KV graph-capturable.
3. Masked grouped GEMM. MoE's data-dependent \(M_i\) would change shapes every step; the fixed-capacity masked layout keeps the shape static while the count varies inside it.

But do not celebrate yet: graphs kill launch overhead; the kernel boundaries remain, and every boundary is a global sync. The expert GEMM cannot start until dispatch has fully finished, and combine cannot start until the GEMM finishes. During communication the SMs idle; during compute the network idles. DeepEP's hooks mitigate this at micro-batch granularity, but overlap can never cross a kernel boundary.

Mega MoE (in DeepGEMM) takes the logical endpoint: fuse the entire MoE layer (dispatch, Linear1, SwiGLU including re-quantization, Linear2, combine) into one kernel, via two enablers. SymmBuffer (symmetric memory): buffers with an identical address layout on every rank of an NVLink domain, so remote GPU memory is read and written with ordinary load/store instructions. Communication becomes a memory instruction inside the compute kernel. Warp specialization, generalized: dispatch warps pull this rank's tokens out of the symmetric buffer and track arrivals; MMA warps run the two GEMMs; epilogue warps do SwiGLU, re-quantization to FP8, and combine write-back, handing off at tile granularity, tokens joining the GEMM as soon as they arrive. No "all dispatch complete" sync point exists. (Precision: FP4 weights + FP8 activations, SM100 primarily, with an SM90 adaptation.)

Which route wins is a hardware question: the fused route's scope is the NVLink-domain size. An 8-GPU domain covers only small EP degrees; NVL72 puts 72 GPUs into one symmetric-memory domain, so decode-scale EP fits entirely inside the fused route's scope. That is the hardware background of its current prominence. And the phase asymmetry returns: decode is IO-bound with a high launch/sync share, so fusion pays off most there; prefill's large GEMMs already run near the compute ceiling.

The 2026 Line: When the Kernel Problem Inverts

The 2026 frontier models changed the kernel workload itself: Kimi K3 and GLM-5.3-Flash make gated delta-rule linear attention (KDA) the dominant layer type, run routed experts in native MXFP4, and draft with block diffusion instead of autoregressive heads. KDA keeps a matrix-valued recurrent state \(S \in \mathbb{R}^{128 \times 128}\) per head (96 heads), updated as \(S_t = (I - \beta k k^{\top})\, \mathrm{Diag}(\alpha)\, S_{t-1} + \beta k v^{\top}\). There is no \(L \times L\) score matrix at all. The kernel problem inverts FlashAttention's: not "avoid materializing \(S, P\)" but "read, update, and write a per-head matrix state efficiently".

On the prefill side the recurrence goes chunkwise: a within-chunk parallel pass plus cross-chunk state passing (chunk size 64, 16-token secondary tiles), the state carried across chunk boundaries exactly like the \((m, \ell, O)\) triple. FlashKDA implements this in CUTLASS for SM90+ (the FLA library auto-dispatches to it). It is the first linear-attention kernel line with FlashAttention-class maturity. At decode, each KDA layer reads its full state, applies one rank-one-ish update, writes it back. Per layer per step: \[ 96\ \text{heads} \times 128 \times 128\ \text{elem} \times 2\ \text{B (BF16)} = \text{3.0 MiB} \quad \text{(fixed, independent of } L\text{)}, \] so K3's 69 KDA layers cost \(69 \times 3.0\) MiB = 207 MiB ≈ 0.2 GiB of state traffic per token-step. And the fused KDA decode kernel does it in one launch (short causal conv of kernel size 4, recurrent update, gated output RMSNorm; all states updated in place, no intermediate tensors, no per-op launches).

MXFP4 MoE in production. K3's routed experts ship native MXFP4 (QAT from SFT onward; 17.5 MB per expert). Executing them is not "FP8 with fewer bits": the SiTU-GLU activation must be fused into the grouped-GEMM epilogue, and Blackwell's native FP4 datapath is what makes 4-bit weights a bandwidth win rather than a dequantization tax. vLLM's recommended K3 backend is deep_gemm_mega_moe (the fused kernel above, in FP4 production), with flashinfer_trtllm for TP > 1.

Common Pitfalls

Research Thinking

How to: read the starting point; read the question and think, for a minute, a day, a week... and only then open the answers. You are not supposed to reinvent several months of someone's research; each paper took its authors exactly that long. It is the habit of thinking that counts.

Have Fun! Squash the Running Max

Enough theory. Now you be the reduction kernel. Below are two KV blocks, each with its local stats. Move the sliders and watch what happens to the old block when the new block's max exceeds it. Can you make the old block's merge weight drop below 0.1? Below 0.01? (Defaults are the first component of our worked example: \(m_1 {=} 2, \ell_1 {=} 8, O_1 {=} 4\) vs \(m^{(2)} {=} 3, \sum e^{s-3} {=} 6, \sum e^{s-3}v {=} 18\).)

Seminar & Homework

Take Lab 7: Kernels (FlashAttention, Launch Overhead, CUDA Graphs) from the hands-on pack. Three experiments, one GPU:

1. Make the traffic argument visible (NGC PyTorch container): one attention layer, 32 heads, \(d_h = 128\), BF16. Time naive (\(QK^\top\) + softmax + \(PV\), materializing \(S, P\)) vs. F.scaled_dot_product_attention at \(L \in \{4K, 16K, 64K\}\) (warm up 20 iterations, time 100 with torch.cuda.Event) and convert to achieved GB/s. The expected picture:

   \(L\)	naive traffic	naive time @ β	flash traffic	flash time @ β	ratio
   4K	3.0 GiB	0.96 ms	128 MiB	0.04 ms	24×
   16K	48 GiB	15.4 ms	512 MiB	0.16 ms	96×
   64K	768 GiB	246 ms	2.0 GiB	0.64 ms	384×

   Both kernels achieve a similar fraction of β, so the traffic ratio, not the FLOPs, explains the gap, and the ratio grows linearly in \(L\) (why? connect it to Lecture 3's \(L^*\) crossover).
2. Count the kernels (vLLM container): nsys profile -o decode_step vllm serve ... at \(B{=}1\), drive 200 single-stream decode requests, open the trace, count kernels between two sampled tokens. Expect ~300+ per 8B decode step (32 layers × 8–12 kernels + glue); at 3–5 µs of CPU launch work each, ~1.0–1.7 ms of launch path against a ~5–12 ms step.
3. Flip the graph switch: repeat with --enforce-eager; compare step time and CPU-side launch gaps. Expect eager to hurt far more at \(B{=}1\) than at \(B{=}64\): fat kernels hide launch latency.

Lab pitfalls, verbatim-worthy: no warm-up (first-call JIT/capture contaminates the trace); timing with time.time() instead of CUDA events (async launches lie); profiling at \(B{=}1\) only and generalizing to serving; counting host time as kernel time. The analysis questions you will answer there: which kernels dominate decode step time (GEMMs, attention, or elementwise), and does it match the \(2P/\beta\) floor?; what exactly do CUDA graphs trade (capture cost, memory, batch-shape rigidity; why does vLLM capture only specific batch sizes)?; at what batch size does --enforce-eager stop mattering, and why?; and where would DeepGEMM-style JIT show up in your trace, and which benchmarking pitfall does that motivate?

Summary

• Kernel design is a data-residency policy over registers / SMEM / L2 / HBM: only the first two are programmer-controlled, and bytes that stay on chip are ~free. Hopper's TMA, WGMMA, and clusters decouple "move bytes" from "multiply bytes": the producer/consumer/epilogue skeleton shared by FA3, DeepGEMM, and Mega MoE.
• Naive attention moves \(4L^2 b_{dtype}\) bytes per query head (\(I \approx d_h/b = 64 < 295\): bandwidth-bound even in prefill, 128 GB per head at \(L{=}128\)K). Online softmax + tiling keeps \(S, P\) on chip: \(O(Ld)\) traffic, 1820× less per layer, and exactly the same math, the \((m, \ell, O)\) recurrence, proved by induction, exact to ~\(10^{-16}\).
• Decode kills Q-parallelism → split-KV with the LSE merge, the same math as DCP at another scale. FlashInfer supplies the engine contract; FlashMLA runs MLA's absorbed shape (128 heads sharing one 576-dim latent: 4.6 GB/step at 128K, ≈ 90% MBU); the DSA chain drops attention traffic \(O(L) \to O(k)\) at the price of two extra kernels.
• Decode MoE has tiny per-expert \(M_i\) (\(\eta_{tile} \approx 6.25\%\)) and pays \(E_{local}\) launches. Grouped GEMM with a masked layout fixes both and stays CUDA-graphable; DeepGEMM adds fine-grained FP8 scaling, two-level accumulation, and JIT; DeepEP runs all-to-all in two phase-specific modes.
• Launch overhead: \(T_{step} = T_{exec} + n_k(t_{launch}+t_{gap})\); ~\(10^3\) kernels × ~3.5 µs ≈ 34% of a 10 ms step. CUDA graphs reclaim it (1.5× at zero kernel cost) given static shapes (buckets at \(1/K\) padding waste, block-table indirection, masked MoE layouts). But boundaries survive as sync points, which Mega MoE erases by fusing comm and compute over symmetric memory, one tile at a time.
• The 2026 line inverts the problem: KDA state costs a fixed 3.0 MiB/layer (0.7% of K3's decode attention traffic at \(L{=}1\)M; the hybrid's speed lives here), MXFP4 experts demand fused-epilogue kernels, and every extra kernel (DSA indexer + top-k, drafters) spends from the launch budget graphs just bought back.