Skip to content

IO-Aware Kernel Design

A recurring GPU-kernel design template, originating with FlashAttention, that treats HBM round-trips of intermediate tensors as the binding bottleneck — not arithmetic intensity. The recipe: identify an N × K intermediate that’s materialized in HBM only to be reduced along K; fuse the produce/consume pair into a streaming kernel with on-chip state; exploit an algebraic decomposition (softmax over blocks, argmax over partitions, segmented sums) that lets the reduction commute with tiling. The result is mathematically exact, not an approximation. Recent filed work generalizes the template off attention onto Lloyd’s k-means and LM-head categorical sampling.

  • [2026-07-16] Kimi K3 — Open Frontier Intelligence (2.8T MoE with KDA + AttnRes): Kimi K3 contributes a KDA-compatible prefix-cache implementation to vLLM — a serving-side kernel dependency for the hybrid-linear-attention backbone — and recommends supernode deployment (≥64 accelerators) so expert-parallel communication stays inside a single high-bandwidth domain.
  • [2026-07-14] Inside TPU and GPU Clusters: The Anatomy of Collective Communication: Documents the multi-GPU-side hardware constraints that the compute-communication kernel work on this page (ThunderKittens ParallelKittens, FA4, SonicMoE, NVSHMEM) is designed to exploit: TPU ICI torus with per-hop latency and bandwidth hierarchy vs GPU fat-tree with NVSwitch SHARP in-network reduction and multicast. Concretely notes SHARP’s ~2× ideal All-Reduce speedup collapses to ~1.3× in practice on NCCL 2.29.7 / H100 IB — the specific gap that motivates kernel-level comm/compute overlap in Loads and Loads of Fluffy Kittens and NVSHMEM device-initiated communication.
  • [2026-07-10] The 4-bitter Lesson: Balancing Stability and Performance in NVFP4 RL: Concrete Blackwell/PTX kernel work for end-to-end NVFP4 RL: (a) per-token FP32 activation scaling fused into the quant kernel to avoid batch-shared scales; (b) a Four-Over-Six fast path using cvt.rn.f16x2.e2m1x2, cvt.rn.f16x2.e4m3x2, and mul.rn.f16x2 to keep FP4·FP8 multiplies lossless in FP16 registers (BF16 lacks mantissa bits), ~2.8× faster than the FP32-scalar reference at >99.97% scale-choice agreement; (c) TransformerEngine PR eliminating duplicate quantized tensor storage, cutting NVFP4 linear peak memory 150 MB → 45 MB (–70%) on a 2k×2k→8k proxy.
  • [2026-06-30] waterloo_intern (ali / Baseten) — Worklog: retrofitting sparsity for the world's fastest video generation kernel (54×): Baseten engineer ali (@waterloo_intern) teases a worklog on retrofitting sparsity into a frontier OSS video generation model at inference time, claiming the world’s fastest video generation kernel and 54× cumulative speedup. Worklog details (sparsity pattern, kernel implementation, baseline) not retrievable at filing time — captured as a tweet-as-pointer until the full writeup is accessible.
  • [2026-06-29] Demystifying NVSHMEM: A System-Level Analysis on Symmetric Memory and Device-Initiated Operations in GPU Communication: Consolidates the NVSHMEM design — symmetric memory + device-initiated put/get + IBGDA — that the compute-communication kernel work on this page (ThunderKittens, FA4, SonicMoE) sits one layer above; identifies where current runtime defaults (completion semantics, fence/quiet cost) leak into upstream kernel performance, complementing Hazy Research’s earlier observation that off-the-shelf NVSHMEM defaults are suboptimal.
  • [2026-05-26] Laguna M.1/XS.2 Technical Report: Laguna fuses MoE all-to-all dispatch/combine into the CUTLASS grouped-GEMM kernel: 8 of 132 SMs per H200 copy expert tokens over NVLink in expert order and set HBM ready-flags, the tile scheduler waits on flags before processing, and the GEMM epilogue stores results directly to remote HBM via shared-memory vectorized stores — concrete production instance of the compute-communication-fusion thesis.
  • [2026-05-24] VideoNSA: Native Sparse Attention Scales Video Understanding: Locates the wall-clock bottleneck inside NSA’s three-branch sparse attention — at 128K video-text context, the compression branch dominates runtime (linear-in-context wall clock), while selection and sliding-window remain cheap. Concrete kernel target for the IO-aware template applied to learnable sparse attention: fuse NSA’s per-block average-pooling + MLP projection with downstream gated combination so the block-level KV tensor is never materialized to HBM. Companion observation (Finding 6) that learnable sparse attention produces branch-specific attention sinks — compression manufactures them, selection eliminates them — adds a numerical-stability dimension worth tracking when designing the fused kernel.
  • [2026-05-24] Attn-QAT: 4-Bit Attention With Quantization-Aware Training: extends the template to training-side co-design with low-precision fused attention. Identifies that FlashAttention’s memory-efficient backward identity (PidPi=dOiOi\mathbf{P}_i^\top \mathbf{dP}_i = \mathbf{dO}_i^\top \mathbf{O}_i) implicitly assumes forward and backward use matching precision; under fake-quantized FP4 forward the identity breaks and gradients become inconsistent. Two fixes: (1) materialize an auxiliary O=PVF\mathbf{O}' = \mathbf{P}\mathbf{V}^F during the forward (25% extra storage) and use it in the backward; (2) fake-quantize the recomputed P\mathbf{P} in the backward to match the forward precision. Result: FP4 attention matches BF16 on Wan-2.1-14B video diffusion (human 2AFC) and Qwen3-14B / Llama-3.1-70B continued pretraining + SFT, without any outlier-mitigation heuristics. Paired CuTeDSL kernel (FlashAttention-4 FP4 on B200) runs block-scaled NVFP4 on QK\mathbf{QK} + BF16 on PV\mathbf{PV} — the “FP4 only where it’s free” design point dictated by Blackwell’s softmax bottleneck — reaching 1801 TFLOPs/s and 1.39× over FA4 BF16 at long context.
  • [2026-05-23] Gram Newton-Schulz: A Fast, Hardware-Aware Newton-Schulz Algorithm for Muon: extends the template to optimizer-state computation. Gram Newton-Schulz rewrites Muon’s Newton-Schulz iteration so the dominant computation runs on the small symmetric Gram matrix XX\mathbf X \mathbf X^\top instead of the rectangular momentum matrix, then exploits custom CuTeDSL symmetric-GEMM kernels (Hopper + Blackwell) that write only the lower triangle. 40–50% wall-clock reduction in the Muon orthogonalization step on trillion-parameter MoEs (Kimi K2-scale), with training quality preserved within 0.01 validation perplexity. Adds “structured-matmul exploitation” as a distinct mechanism alongside HBM-traffic elimination, and introduces a half-precision stability investigation (spurious negative eigenvalues amplified by (15/8)2(15/8)^2 per step, plus eigenvector drift) fixed by periodic restart.
  • [2026-05-23] FlashSampling: Fast and Memory-Efficient Exact Sampling: extends the template to LM-head categorical sampling. Argmax-over-partitions plays the same role for sampling that softmax-over-blocks plays for attention; the categorical distribution’s hierarchical factorization enables a tensor-parallel variant that replaces all_gather(logits) with streaming per-rank (max, argmax) exchanges.
  • [2026-05-23] Flash-KMeans: Fast and Memory-Efficient Exact K-Means: applies the template to Lloyd’s k-means. FlashAssign fuses streaming distance computation with online argmin (no N×K distance matrix in HBM); Sort-Inverse Update converts scatter-atomic centroid updates into argsort + segmented partial sums. Up to 200× over FAISS at iteration latency.
  • [2026-05-23] FlashAttention-4: Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling: re-tunes the original FlashAttention design for Blackwell. New roofline accounts for SMEM traffic and MUFU exp throughput as co-dominant with matmul; introduces software-emulated 2^x on FMA units, conditional softmax rescaling, and a backward pass that uses 2-CTA MMA to halve operand-B SMEM staging.
  • Where else does the template apply on the wiki? Any kernel that materializes an N × K intermediate solely to reduce along one axis is a candidate: top-k logits, MoE router gating, perceptual losses, RoPE applied before attention, contrastive logits. None of these have been filed yet as IO-aware rewrites. Partial answer: Gram Newton-Schulz is an adjacent variant — instead of fusing a produce/reduce pair, it rewrites the iteration so the dominant matmul is structured, and ships a kernel that exploits the structure. This suggests an enlarged template: “find an intermediate whose algebraic structure (symmetry, low rank, banded) is implicit but not exploited by standard libraries, restructure the algorithm to make the structure dominant, write the structured kernel.” VideoNSA adds a learnable sparse candidate: NSA’s compression branch is N × K (per-block mean-pool + MLP) → reduced by the per-head gate; fusing the pool + projection + gate combination is the obvious next target.
  • Backward-pass coverage. FA4 has a backward pass; Flash-KMeans, FlashSampling, and Gram Newton-Schulz are inference / online primitives only. Attn-QAT is the first filed paper that targets the backward pass head-on, showing that the FlashAttention backward identity implicitly assumes forward/backward precision match — and that under low-precision training this assumption must be made explicit by carrying a small high-precision auxiliary output. Whether the template applies to backward passes that don’t have an obvious partition decomposition (e.g. CE loss + LM-head backward, which is exactly the locus that Lost in Backpropagation: The LM Head is a Gradient Bottleneck flags as a bottleneck) is still unaddressed.
  • Compile-time as the bottleneck. All filed papers either rebuild compile infrastructure (FA4’s, Gram-NS’s, and Attn-QAT’s CuTeDSL moves) or replace autotune sweeps with analytical heuristics (Flash-KMeans’s cache-aware config selector). The implicit claim is that iteration speed on the kernel-design loop itself is now load-bearing for architecture-specific tuning at the rate Blackwell/Hopper/B300 ship; whether this means Python-DSL becomes the dominant paradigm or whether C++-template + cache modeling wins is open. Dao-AILab’s stack (FA4 + Gram-NS + quack/gemm_symmetric + Attn-QAT FA4-FP4) is the strongest existing endorsement of the CuTeDSL-in-Python direction.
  • Approximation-as-extension. Three of the original four papers emphasize exactness. The natural next axis is controlled approximation — block-sparse attention, approximate k-means with bounded error, top-k sampling — built on the same kernel skeleton. The wiki has MonarchRT: Efficient Attention for Real-Time Video Generation as one such datapoint on the attention side; Attn-QAT is the precision-side version of approximation — quantization is a controlled approximation, and the training-time fix lets it match BF16 in practice. VideoNSA / NSA is the learnable sparse version of the same axis: sparsity pattern + gating are trained, not statically chosen.
  • Numerical stability in half precision. Gram Newton-Schulz’s restart trick is the explicit instance of this: bfloat16 spurious negatives blow up exponentially via rt<(15/8)2rt1r_t < (15/8)^2 r_{t-1}. FA4’s conditional softmax rescaling addresses a milder version of the same problem. Attn-QAT adds a third instance: the FlashAttention backward identity silently breaks under precision mismatch, and only the auxiliary O\mathbf{O}' trick + matched backward recomputation restores it. Whether other applications of the template (Flash-KMeans, FlashSampling) have undiagnosed half-precision pathologies is open.
  • Where does the FP4 attention story go next? Attn-QAT’s design point on B200 — NVFP4 QK\mathbf{QK} + BF16 PV\mathbf{PV} — is dictated by the softmax bottleneck; B300 doubles MUFU exp throughput and Rubin quadruples it, which should make NVFP4/MXFP8 QK\mathbf{QK} + FP8 PV\mathbf{PV} a net win. The open meta-question is whether the QAT recipe needs to be re-derived per hardware generation, or whether the precision-consistency fixes (auxiliary output + matched backward recomputation) transfer.