Topic
ML & GPU
Deep dives into machine learning systems, GPU programming, LLM inference, attention mechanisms, speculative decoding, vector search, and model optimization. From CUDA kernels to production serving.
75 posts · ~551 min of reading
Oct 16, 2026
Quake's Recall Estimator Works Because Its Geometry Is Off: Caps in d/2, a 1/e Ceiling, and a Split Example That Needs α < 0.51
Quake (OSDI '25) chooses nprobe per query by estimating recall from hyperspherical cap volumes, with no offline tuning. On SIFT1M I reproduce its Table 5 to within a point of recall. The released code computes cap volumes with sqrt(1−t²) where the textbook formula has 1−t², which is exactly the volume in 64 dimensions instead of 128. That error is what makes the estimator calibrated: the textbook formula delivers 78% recall at a 90% target. The same audit finds a 1/e ceiling on the first-partition estimate and a split example that only commits when α < 0.509. Measured α is 0.67–0.83.
Oct 12, 2026
Do You Need the 'H' in HNSW? Flat Graphs, Hub Highways, and Where the Distance Computations Actually Go
An ICML 2025 workshop paper finds that a flat navigable small-world graph matches HNSW on 13 high-dimensional benchmarks, and argues that hub nodes form a 'highway' that does the hierarchy's job. In my own 50k-point simulation the flat graph matches HNSW to within 2% at every d ≥ 8, including on the uniform sphere, which has no hubs at all. The upper layers account for just 6–9% of distance computations there, so the hierarchy has little left to save.
Oct 12, 2026
SuperBPE on a 1.3 MB Corpus: Superword Tokens Cut 21% of Tokens In-Domain, Half That Out of Domain, and My Section Headings Became Vocabulary
SuperBPE (COLM 2025) trains BPE in two stages: subwords first, then merges that are allowed to cross whitespace. I reimplemented it and trained it on this blog's own posts. At a 16k vocabulary it needs 20.9% fewer tokens than BPE on held-out posts, but only 10.9% fewer on npm READMEs. The transition point that is best in-domain is close to the worst out of domain. The paper's 4-word cap had no effect on encoding efficiency, and it also failed to stop the tokenizer from learning my post template ('.\n\n## What the') and markdown table rules as tokens.
Oct 10, 2026
SynthID-Text's Tournament Watermark: 2^30 Candidates Become One Line per Layer, and the 10→1 Weights Are 99% Optimal
SynthID-Text (Nature, 2024) watermarks LLM output with a 30-layer knockout tournament over sampled tokens, but the reference code never samples a candidate. Each layer is the closed-form tilt p·(1 + g − G). I verified the tilt against a literal tournament, derived how it concentrates the distribution, and ran the watermark end to end on Qwen2.5-0.5B. Per-layer signal falls 6× from layer 1 to layer 30. The default 10→1 weights capture 99.1% of the optimal signal-to-noise ratio, and the plain mean needs 25% more tokens.
Oct 8, 2026
Cut Cross-Entropy: Training Loss Without the Logit Matrix, and Why 0.1% Element Sparsity Becomes 50% Block Sparsity
Cut Cross-Entropy (Apple, ICLR 2025) computes the LM loss and its gradient without ever writing the N×|V| logit matrix to HBM, cutting Gemma 2 2B's loss memory from 24 GB to 1 MB. Its 1,161 MB 'lower bound' row is exactly the size of the bf16 gradients for C and E. The speed comes from skipping softmax blocks below 2⁻¹². I measured that on Qwen2.5-0.5B: 0.125% of softmax entries clear the threshold, yet 49% of 128×128 blocks still have to be computed (99% under a random vocab order). At the paper's sparsity level my model keeps 21% of blocks, which matches the ~20% implied by the paper's own timings. Filtering also zeroes the classifier gradient for 10% of the vocabulary, which is why pretraining needs the FullC variant.
Oct 8, 2026
FlashAttention-4's Softmax Tricks: A Degree-3 exp2 on the FMA Pipe, and a Lagging Max That Saturated FP8
On a B200 the tensor cores do 8192 FLOPs per clock per SM, while the exp2 unit does 16. FlashAttention-4 (arXiv:2603.05451) shifts some exponentials onto FMA units with a degree-3 polynomial, and skips output rescaling unless the row max grows by more than 2⁸. I rebuilt both in NumPy. The polynomial matches the paper's Table 2 to three digits. It is bit-identical to correctly rounded bf16 on 99.0% of inputs and within 1 ULP on 100%. Without the threshold, 6–95% of row-blocks rescale in my runs; with it, 0–16% do. But its exactness argument assumes P can hold up to 2⁸ without loss. An FP8 E4M3 P can't, and FA4's own issue #2716 hit exactly that.
Oct 8, 2026
QTIP: Trellis-Coded Quantization Puts 2-Bit LLM Weights 10% Above the Rate-Distortion Bound, and a Monotone Codebook Loses to Scalar Rounding
Vector quantizers for LLM weights stop at 8 dimensions because the codebook has to fit in cache. QTIP (NeurIPS 2024) uses trellis-coded quantization instead: a 256-dimensional quantizer whose decoder is a bit shift plus three ALU instructions. I wrote a bitshift-trellis Viterbi quantizer from scratch. It reproduces the paper's 2-bit Gaussian MSEs to within 0.001 (0.0697 vs 0.069, and 0.0723 vs 0.0733 for the tail-biting table). It also shows the failure the paper only plots: a codebook that is monotone in the state index gets 0.143 MSE, worse than the 0.118 of a plain 4-level scalar quantizer, at every trellis size.
Oct 7, 2026
mHC: Sinkhorn-Projected Residual Streams, Where the 20-Iteration Error Actually Lands, and Why Fig. 8's 1.50 Column Sum Is a Correlation Effect
DeepSeek's mHC widens the residual stream to n=4 lanes and projects each lane-mixing matrix onto the Birkhoff polytope with 20 Sinkhorn iterations. The paper reports a peak composite gain of 3000 under plain HC, and 1.6 under mHC. I simulated the projection. Because row normalization runs last, the forward gain is exact to 1e-15 and all the leftover error goes to the backward (column) side. Independent random layers only drift to 1.14. Layers with shared structure reproduce the paper's 1.50 / 0.41 column sums. The paper's max-only metric hides the 0.41.
Oct 5, 2026
TurboQuant: Random Rotation, Lloyd-Max Codebooks, and the QJL Residual Trick for KV-Cache Quantization
TurboQuant (Google Research, arXiv:2504.19874) quantizes vectors online with no calibration: rotate randomly, then apply one fixed Lloyd-Max scalar codebook to every coordinate. Its MSE comes within 1.44–2.7× of the Shannon lower bound. I reimplemented it at head dimension 128. The MSE numbers reproduce exactly. The b=4 inner-product constant in Theorem 2 is understated by 12%. And the paper's unbiased 'prod' variant loses to a simpler fix, multiplying by one scalar, at every bit width I tested. In a needle-in-softmax test it had up to 2× the attention-output error.
Oct 4, 2026
Gated Attention: A Sigmoid After SDPA, the Low-Rank W_V·W_O Bottleneck, and Attention-Sink-Free LLMs
The Qwen team's NeurIPS 2025 best paper found that one head-specific sigmoid gate on the SDPA output lowers perplexity, removes loss spikes, and cuts first-token attention from 46.7% to 4.8%. I went through the paper's tables. Across six gated variants, the mean gate score ranks sink mass and massive activations exactly (Spearman 1.0, n=6). The RULER gain is mostly a context-extension effect: +0.3 to +1.7 points inside the trained length, +35 points at 32k after YaRN. A simple length model shows why a learned sink can't stay calibrated when the context grows.
Sep 29, 2026
FlexAttention: Ending the Attention Software Lottery with Compiler-Generated Kernels
FlexAttention (Dong, Feng, Guessous, Liang & He, arXiv 2412.05496) turns attention variants from a kernel-engineering problem into a four-argument Python function: torch.compile inlines a user-written score_mod/mask_mod into handwritten Triton templates and derives block sparsity from a BlockMask. The result covers alibi, sliding window, PagedAttention and their combinations at 0.68x–1.43x of FlashAttention-2 — and 5.37x faster where FlashAttention's fallback loses the lottery. I reproduced the BlockMask classification locally: causal at 4k computes 528/1024 blocks, 93.9% of them mask-free, and the mask shrinks from 16 MiB dense to 4.1 KiB.
Sep 28, 2026
ThunderKittens: 16x16 Tiles, the Load-Compute-Store-Finish Template, and Why 40 Lines Can Match cuBLAS
ThunderKittens (Spector et al., arXiv 2410.20399) bets that three abstractions — 16x16 register/shared tiles with compile-time swizzled layouts, a producer-consumer thread-block template, and a persistent grid with cache-aware block ordering — cover almost everything an H100 AI kernel needs. The receipts: a 40-line GEMM that matches cuBLAS, attention backward 10-40% faster than FlashAttention-3, and 8-14x speedups on linear attention and SSMs. I re-derived the bank-conflict claim at the center of its layout system with a 30-line simulator: naive row-major really is 8-way conflicted, and the 128-byte XOR swizzle really is conflict-free.
Sep 28, 2026
WaferLLM: The PLMR Model, Interleaved Cannon Shifts, and LLM Inference on 850,000 Cores
WaferLLM (OSDI '25) is the first LLM inference system built for wafer-scale chips like the Cerebras WSE-2: 850,000 cores, 48 KB of SRAM each, arranged in a 2D mesh with no shared memory. The paper's PLMR device model explains why GPU-style allgather, SUMMA, and even Cannon's algorithm all break on this hardware, and the fixes — an interleaved cyclic shift that bounds every communication step to two physical hops, a K-tree allreduce for GEMV, and shift-based KV cache placement — deliver 606x faster GEMV than an A100 and 10-20x end-to-end speedups over multi-GPU SGLang. The constraints are as interesting as the wins.
Sep 27, 2026
DeepSeek Sparse Attention: A Lightning Indexer, Top-k Token Selection, and Retrofitting Sparsity onto a Trained Model
DeepSeek-V3.2-Exp swaps dense attention for fine-grained sparsity — not by training from scratch, but by bolting a tiny FP8 'lightning indexer' onto V3.1-Terminus and distilling the model's own attention distribution into it via KL loss. Each query attends to only the top-2048 of 128K tokens (1.6% density), core attention drops from O(L²) to O(Lk), and decode cost at 128K goes nearly flat. The interesting part is the training recipe: 2.1B tokens of frozen-model warm-up, then 943.7B tokens of sparse adaptation with the indexer detached from the main gradient graph.
Sep 23, 2026
MUVERA: Collapsing Multi-Vector Retrieval into One MIPS Query
ColBERT-style late interaction scores a query against a document with Chamfer similarity over token embeddings — a sum of maxes that no off-the-shelf vector index can serve. MUVERA (Google, NeurIPS 2024) compresses each token set into one Fixed Dimensional Encoding whose plain inner product ε-approximates Chamfer, turning the whole problem into standard single-vector MIPS: 10% higher recall than PLAID at 90% lower latency, with a provable guarantee. A from-scratch reimplementation shows why it works — the estimator is heavily biased (−0.21) but almost uniformly so (σ=0.02), and the asymmetric sum/average construction plus fill_empty_clusters carry the ranking.
Sep 15, 2026
Aegaeon: Token-Level Auto-Scaling for GPU Pooling Across Many LLMs
Aegaeon (SOSP '25) preempts models between decode steps instead of between requests, packing up to seven LLMs per GPU. A queueing argument shows why request-level auto-scaling caps out below three models per GPU, and a full-stack teardown of engine re-initialization — component reuse, bump-allocated VRAM, slab-allocated KV caches, CUDA-event-synchronized swaps — cuts preemptive scaling from 26.9 seconds to under one. Production deployment dropped a serving fleet from 1,192 GPUs to 213.
Sep 14, 2026
Ozaki Scheme II: Emulating FP64 Matrix Multiply with Residue Arithmetic on Integer Tensor Cores
GPUs are quietly abandoning native FP64, and the replacement is number theory. Ozaki Scheme II (IJHPCA 2026) emulates double-precision GEMM by mapping scaled integers into a residue number system, running one exact INT8 tensor-core GEMM per modulus, and reconstructing via the Chinese Remainder Theorem — linear GEMM count where classic mantissa slicing is quadratic. My numpy prototype reproduces the mechanics: 8 residue GEMMs recover a 2^56-magnitude integer product bit-exactly, and the accuracy dial is just the modulus count.
Sep 11, 2026
AutoUVM: Tensor-Level Prefetching When Your LLM Doesn't Fit in GPU Memory
CUDA Unified Virtual Memory lets an oversized model run on a small GPU — at fault-driven migration speeds that make it nearly useless. AutoUVM (arXiv:2609.06172) closes the semantic gap between PyTorch and the UVM driver: it learns the kernel-to-tensor access sequence offline, classifies each kernel against a three-roof roofline, and issues cudaMemPrefetchAsync only where the CPU–GPU interconnect is the binding roof. Result: 3.1x over stock UVM at 2x oversubscription. Inverting a two-roof cost model against the paper's own sensitivity numbers implies stock UVM's redundant migration grows superlinearly — roughly 1.1x, 2.2x, 2.8x the minimum overflow bytes at 1.5x, 1.75x, 2x oversubscription.
Sep 10, 2026
Marconi: Prefix Caching When Your Model Can't Roll Back Its Own State
Hybrid SSM/attention models break the core assumption of prefix caching: recurrent states update in place, so a cached state for tokens 1..q is useless for a request that shares only 1..p. Marconi (MLSys 2025 Best Paper) fixes this with speculative radix-tree admission and FLOP-aware eviction, lifting token hit rates up to 34.4x over per-block checkpointing. My own eviction sim reproduces the win on an agentic trace (+10 hit-rate points) — and shows the same policy halving the hit rate on a chat trace, which is exactly why Marconi auto-tunes its eviction knob.
Sep 9, 2026
Lance: Random Access in Columnar Storage Without Giving Up Scans
Columnar formats are supposedly bad at point lookups. Lance (arXiv 2504.15247) shows the real culprit is how formats encode structure — validity and repetition — not columnar layout itself. Two adaptive encodings, mini-block and full-zip, cap any random access at 2 IOPS regardless of nesting depth. The numbers check out: default Parquet manages ~5.5K rows/sec where tuned layouts hit 350K, a 63.6x gap my own arithmetic reproduces.
Sep 1, 2026
Prism: Symbolic Superoptimization, or How to Not Decide Things Yet
Tensor-program superoptimizers die on a product search space: graph structures × tensor-partition mappings × tile sizes. Prism symbolizes the last two factors so the generator never enumerates them. I rebuilt the mapping space by brute force to see where the win actually comes from — and it is a very large constant, not a smaller exponent.
Aug 31, 2026
The Checkpoint Tax at 16k GPUs: Daly's Formula Stopped Working
Every systems course teaches the Young/Daly optimal checkpoint interval. A wave of 2025-2026 LLM training papers quietly ignores it and checkpoints far more often than the formula allows. I simulated why, then verified the trick that makes it safe — torn snapshots repaired by replaying the optimizer on the host — and found a 1.9x traffic reduction the papers leave on the table.
Aug 27, 2026
Fair Queueing for LLM Serving: Why the Virtual Token Counter Needs a Lift Rule
Requests-per-minute limits are the industry's fairness mechanism, and they are terrible: to get a tighter fairness gap than a proper fair scheduler, my simulation had to throw away 78% of the GPU. VTC (OSDI '24) ports weighted fair queueing to continuous batching, and the part that matters is not the counter — it is the one line that erases a returning client's banked credit. Without it, an idle client comes back and takes 83.5% of the machine.
Aug 23, 2026
Microscaling FP4 Attention: Why the Softmax Operand Needs Two Scale Factors
SageAttention3 runs both attention matmuls in NVFP4 on Blackwell tensor cores at 1038 TOPS. The interesting part isn't the format — it's that post-softmax probabilities live in [0,1], which collapses an E4M3 scale factor to 35 usable codes. A from-scratch reimplementation shows the two-level fix is worth 30 cosine-similarity points, but only for high-entropy attention.
Aug 22, 2026
Tiling Is Function Composition: The Partial Algebra Behind Every Tensor Core Kernel
CUTLASS turns tiling, swizzling and thread assignment into a single algebraic operation on shape/stride tuples. I implemented the algebra from scratch to check it, and found the interesting part is where it breaks: composition is a partial operation, and I measured how partial. For a 4x4 row-major layout only 76% of candidate composands are legal — which is exactly why these conditions are static_asserts.
Aug 19, 2026
One Nonlinearity Left: Cross-Layer Transcoders and the Mechanics of Attribution Graphs
Sparse autoencoders give you a dictionary of features but no mechanism. Cross-layer transcoders plus frozen attention collapse a transformer into a network with exactly one nonlinearity left, which makes a circuit an exact linear decomposition rather than a gradient approximation. I rebuilt the construction in numpy, confirmed the decomposition is exact to 1e-15, and localized where its faithfulness actually breaks.
Aug 18, 2026
Worst-Case Optimal Joins on a GPU: Flattening Skew Into a Prefix Sum
Worst-case optimal joins fix the asymptotic blowup that kills binary join plans on cyclic queries, but their attribute-at-a-time intersections are a disaster on SIMT hardware: one high-degree hub key and a single warp holds up the whole device. SRDatalog's answer is to stop scheduling keys and start scheduling work units, flattening a skewed multi-level search space into a prefix-summed 1-D array that thread blocks slice at kernel launch. I re-derived the AGM bounds behind their rule rewriting and simulated their scheduler, and the numbers say the residual bottleneck is not where the paper says it is.
Aug 16, 2026
Tensor Memory: Why Blackwell's MMA Is Issued by One Thread
Blackwell's tcgen05 instructions break two rules that every GPU matrix-multiply kernel has been built around: accumulators no longer live in registers, and the MMA is no longer collective. I worked through the PTX ISA's Tensor Memory chapter and derived the arithmetic that forced both changes — accumulator traffic per FLOP scales as K/4, and every new tcgen05 kind lands at exactly 2x Hopper's per-SM accumulator bandwidth. That number is the whole design.
Aug 15, 2026
vAttention: Dynamic KV Cache Allocation Without Breaking Virtual Contiguity
PagedAttention solved KV cache fragmentation by making the cache non-contiguous in virtual memory, and every attention kernel since has paid for that decision in rewrites, register spills, and block-table bookkeeping. vAttention points out that the fragmentation was always a physical-memory problem, and fixes it with CUDA's virtual memory APIs instead. I re-derived the paper's page-size tables and its batch-size result from first principles; both hold, and the arithmetic reveals why they had to patch NVIDIA's driver to ship it.
Aug 14, 2026
Running TPC-H on NCCL: The Broadcast That Refuses to Scale
A distributed SQL engine built on the collective communication library used for ML training runs all 22 TPC-H queries at 1TB in 0.53 seconds on 40 H100s. The abstract leads with that. Their own breakdown figure leads somewhere else: 8 GPUs in one machine do it in 1.13 seconds, so 5x the hardware bought 2.13x. This post derives why, verifies the paper's shuffle-vs-broadcast crossover condition, finds a non-monotonicity in it the paper doesn't mention, and computes the cost-per-query that makes scaling out the worse deal.
Aug 13, 2026
Semantic Operators: Putting a Statistical Contract on 6 Million LLM Calls
A semantic join over 250 documents and 24,370 labels costs 6,092,500 LLM calls and 25 days. Getting that to 5,290 calls is the easy part. The interesting part is what the accuracy guarantee actually covers, and what it quietly does not.
Aug 10, 2026
Kernel-Initiated RDMA: Why Large-Scale MoE Decode Broke the CPU Proxy
Every GPU that sends a network message has to get a work request into a NIC queue, and for years a CPU thread did that on the GPU's behalf. Large-scale expert parallelism broke that arrangement, but not for the reason usually given. The proxy round trip is only 2.3% of a decode layer's budget. This post derives the number that actually matters, checks DeepEP's published latency table against its own bandwidth column, and finds that one of the DeepSeek-V3 hardware paper's design justifications compares against a baseline that never occurs.
Aug 7, 2026
MoE Decoding Is Two Different Workloads: Attention-FFN Disaggregation and the Ping-Pong Pipeline
Sparse activation turns MoE feed-forward layers from compute-bound into memory-bound, and no amount of batching fixes it while attention and FFN share a GPU. MegaScale-Infer splits them onto separate nodes and shuttles micro-batches between them. I re-derived its roofline math, built my own balance model, and found the paper's heterogeneous hardware choice is provable from the price sheet alone, while one sentence in its ablation contradicts its own constraint.
Aug 6, 2026
Processing-in-Memory Has the Bandwidth. It Was Wasting 85 Percent of It on Command Stalls
DRAM-PIM puts MAC units inside every bank, so attention decode should stop being memory-bound. Measured MAC utilization at a head dimension of 128 is 14.7 percent. PIMphony (HPCA 2026) shows all three causes live in the control path, not the datapath: head-first channel partitioning that starves channels, static command scheduling that serializes I/O against compute, and physical addresses baked into precompiled instructions that force worst-case KV allocation. Fixing them costs under 5 percent area and yields up to 11.3x.
Aug 5, 2026
CacheBlend: Reusing KV Caches When Your Text Isn't a Prefix
Prefix caching only works when reused text sits at position zero. CacheBlend reuses KV caches from arbitrary positions and repairs the missing cross-attention by recomputing just 15% of tokens, chosen by a layer-cascading greedy filter.
Aug 5, 2026
FlashInfer: Every KV Cache Layout Is Just a Block-Sparse Matrix
Paged, ragged, shared-prefix, tree-decoded, and pruned KV caches look like five different kernels. FlashInfer collapses them into one block-sparse format with a tunable block size, then adds a load-balancing scheduler that survives CUDAGraph capture.
Aug 4, 2026
Asynchronous RL for Reasoning Models: How Decoupling Rollout from Training Breaks PPO's Math
Synchronous RL post-training wastes most of its inference fleet waiting on the single longest reasoning trace in each batch. Decoupling generation from training recovers that idle time, but it quietly invalidates the assumption PPO's importance ratio is built on, and naive async training drops AIME24 accuracy from 42.0 to 23.3. AReaL fixes the schedule and the objective together.
Aug 3, 2026
TokenWeave: Why Hiding Tensor-Parallel Communication Needs Wave Arithmetic, Not Finer Tiles
Tensor-parallel LLM inference burns up to 20% of its latency in AllReduce, and every framework ships with compute-communication overlap turned off by default. The reason is not laziness: splitting work finer to create overlap costs more than the communication it hides. TokenWeave (MLSys 2026) fixes this with two unglamorous ideas, an unequal split sized to GPU wave boundaries and a fused AllReduce-RMSNorm kernel that runs on 8 SMs out of 132.
Jul 29, 2026
Batch-Invariant Kernels: Why LLM Inference Is Nondeterministic and How to Fix It
Temperature-zero LLM endpoints still return different completions run to run. The usual explanation, floating point plus concurrency, is wrong. The real cause is that reduction kernels change their arithmetic order with batch size, and fixing it turns RL policy divergence from 0.001 into exactly zero.
Jul 29, 2026
Seismic: Why WAND Fails on Learned Sparse Embeddings, and What Replaces It
Learned sparse models like SPLADE produce vectors that look like inverted-index input but behave nothing like it. Exact retrieval with a state-of-the-art WAND implementation takes 100 milliseconds per query on MS MARCO, worse than brute-force dense search. Seismic fixes this by throwing out document-ID ordering entirely: statically prune each list, cluster it into geometrically cohesive blocks, attach a quantized upper-bound summary to each block, and skip blocks whose summary cannot beat the current heap. The result is 187 microseconds at 90% recall, single threaded.
Jul 28, 2026
GPU Confidential Computing: Why the Same Hardware Costs 7% or 41x
Two measurement studies on Hopper-class GPU TEEs report overheads of under 7% and up to 41.6x. Both are correct. The entire gap is explained by one number in NVIDIA's own documentation: the encrypted CPU-GPU path moves about 4 GB/s, against a device that does roughly 990 TFLOP/s. That ratio gives you a break-even arithmetic intensity, and it tells you in advance which workloads survive encryption and which ones fall off a cliff.
Jul 27, 2026
ACORN: Why Filtered Vector Search Breaks HNSW, and the Two-Hop Fix
Every production vector search is really a filtered vector search: find similar documents, but only for this tenant, only from last quarter, only in these languages. HNSW handles this badly, and not for the reason most people assume. The problem is that HNSW's edge pruning rule is provably wrong once you restrict the graph to a subset of nodes. ACORN fixes it by making pruning predicate-agnostic and recovering the deleted edges at query time through two-hop expansion, reaching 2x to 1,000x higher throughput at fixed recall.
Jul 27, 2026
Helix Parallelism: Why Tensor Parallelism Hits a Wall at Million-Token Context
Tensor parallelism has a hard ceiling that most people never hit, because you only hit it when your KV cache is measured in millions of tokens. Once TP width exceeds the number of KV heads, adding GPUs stops reducing per-GPU KV traffic entirely, and the arithmetic says so plainly. Helix Parallelism breaks the ceiling by using a different sharding strategy for attention than for the FFN, in the same layer, on the same GPUs, microseconds apart.
Jul 26, 2026
Megakernels: Deleting the Kernel Boundary to Win Back Microseconds
A single-token forward pass through a 1B parameter model should take 740 microseconds on an H100. Production engines take two to three times that, and the reason is not arithmetic, not memory bandwidth, and not the model. It is the roughly one hundred kernel launches per forward pass, each of which drains the memory pipeline dry. Two 2025 projects independently arrived at the same fix: collapse the entire model into one kernel launch and build a scheduler inside it.
Jul 26, 2026
NanoFlow: LLM Serving Is Compute Bound, and Your GPU Is Half Idle
Everyone repeats that LLM inference is memory bound. Run the arithmetic on a batched 70B serving workload and the opposite falls out: compute is the binding constraint, and yet aggregate GPU utilization sits near 40 percent. The gap is not kernel quality, it is that compute, memory, and NVLink operations execute one after another inside a single device. NanoFlow splits each batch into nano-batches so those three resources run at once, and lands 1.91x over TensorRT-LLM.
Jul 25, 2026
Gated DeltaNet: Householder Transitions and the Arithmetic of Forgetting
Mamba2 forgets everything a little; DeltaNet forgets one thing precisely. Gated DeltaNet does both by replacing the scalar decay in a linear RNN with a scaled Householder matrix, then recovers tensor-core throughput with a WY representation that turns 64 sequential rank-1 updates into three matmuls.
Jul 24, 2026
Lookahead Decoding: Breaking the Sequential Barrier Without a Draft Model
Autoregressive decoding forces one token per forward pass, leaving GPUs starved for work. Lookahead decoding reframes greedy generation as solving a nonlinear system by Jacobi iteration, mining n-grams from the iteration trajectory to collapse many steps into one, with no draft model and no training.
Jul 24, 2026
Microscaling Formats: Why One Shared Exponent per 32 Numbers Changes 4-Bit Training
A single FP32 scale per tensor cannot survive 4-bit quantization, and a scale per element costs as much as the data. Microscaling (MX) formats split the difference — 32 elements share one 8-bit power-of-2 exponent — and that is now silicon in NVIDIA Blackwell tensor cores. The subtle part is not the format. It is that rounding the shared exponent the obvious way silently diverges an 8-billion-parameter pre-training run, and the fix is one direction of a rounding rule.
Jul 20, 2026
Diffusion Language Models: Generating Text Without Left-to-Right
Every large language model you have used generates one token at a time, left to right. LLaDA throws that away: it masks a whole sequence, then denoises it in parallel over a handful of steps — a diffusion process over discrete tokens that matches LLaMA3 8B, beats it on math, and quietly solves the reversal curse GPT-4o still fails.
Jul 20, 2026
Multi-Token Prediction: Training LLMs to See Further Than One Token
Next-token prediction is a strangely myopic objective: at every position the model is graded on exactly one token and told nothing about the future it is steering toward. Multi-token prediction asks the model to forecast several tokens at once. It costs almost nothing at training time, makes larger models measurably better at code and reasoning, and hands you a 3x inference speedup for free through self-speculative decoding — which is why DeepSeek-V3 bakes it into pretraining.
Jul 19, 2026
Titans: Neural Memory That Learns to Memorize at Test Time
Attention is precise but quadratic; recurrent state is cheap but forgetful. Titans adds a third component — a neural memory module that runs gradient descent on itself during inference, storing what surprises it and forgetting the rest — and scales past 2M-token context windows.
Jul 18, 2026
Differential Transformer: Noise-Canceling Attention by Subtraction
Softmax attention leaks probability mass onto irrelevant tokens, a floor of noise that grows with context length. The Differential Transformer borrows a trick from analog electronics — subtract two attention maps to cancel the common-mode noise — and matches a standard Transformer using roughly a third fewer parameters.
Jul 18, 2026
QuaRot: Rotating Away the Outliers That Make 4-Bit LLMs Impossible
The reason you cannot naively quantize an LLM to 4 bits is a handful of activation channels with values 100x larger than the rest. QuaRot (2024) makes those outliers disappear by multiplying the network with random orthogonal rotations that leave the output identical but flatten the value distribution. Here is why an orthogonal matrix is free, why Hadamard matrices make it fast, and how the whole forward pass ends up in INT4.
Jul 17, 2026
Constrained Decoding: How to Force an LLM to Speak Valid JSON Without Slowing It Down
Prompting a model to 'return valid JSON' is a wish, not a guarantee. Constrained decoding turns it into an invariant by masking logits against a finite-state machine or pushdown automaton. The hard part is doing it without adding latency to every single token, which is where compiled indices and context-token classification come in.
Jul 14, 2026
Zero Bubble Pipeline Parallelism: Splitting the Backward Pass to Fill the Gaps
Pipeline parallelism wastes GPU time in bubbles, idle gaps at the warmup and cooldown of every training step. By splitting the backward pass into its input-gradient and weight-gradient halves and deferring the optimizer's sync, ZB-H1/H2 schedules drive the bubble toward zero, buying up to 31% more throughput under synchronous semantics.
Jul 12, 2026
Muon: Orthogonalized Momentum and Why It Trains LLMs 2x Cheaper
AdamW treats every weight as a bag of scalars. Muon treats a weight matrix as a matrix, orthogonalizing the momentum update with a bf16 Newton-Schulz iteration. The result: roughly 2x the compute efficiency at scale, proven on a 16B-parameter MoE trained on 5.7T tokens.
Jul 9, 2026
Contextual Sparsity: Activation-Aware LLM Inference and the Deja Vu Paradigm
Modern LLMs waste 95% of computation on neurons that produce near-zero activations. Contextual sparsity exploits input-dependent activation patterns to skip irrelevant neurons at inference time, achieving 2-6x speedups with negligible quality loss.
Jul 9, 2026
DiskANN: Billion-Scale Vector Search on a Single Machine with Vamana Graphs
How Microsoft Research's DiskANN system uses the Vamana graph algorithm, SSD-optimized layout, and PQ-based beam search to serve billion-scale approximate nearest neighbor queries from a single commodity machine, eliminating the need for distributed in-memory indices.
Jul 9, 2026
Mixture-of-Experts: Expert Parallelism, All-to-All Routing, and Auxiliary-Loss-Free Load Balancing
How MoE models like Mixtral and DeepSeek-V3 route tokens to sparse expert networks across GPU clusters, and why auxiliary-loss-free routing solves the capacity collapse problem without degrading model quality.
Jul 9, 2026
Multi-Head Latent Attention: Compressing KV Cache by 93% Without Losing Quality
DeepSeek-V2 introduced Multi-head Latent Attention (MLA), which replaces per-head KV storage with a shared low-rank latent vector. By jointly compressing keys and values into a bottleneck representation and absorbing the up-projection into attention weights, MLA cuts KV cache memory by over 93% while matching or exceeding standard multi-head attention quality. Here is how the linear algebra works and why this changes the economics of long-context serving.
Jul 9, 2026
PagedAttention: Virtual Memory for LLM KV Caches
How borrowing virtual memory concepts from operating systems, specifically non-contiguous paging and demand allocation, eliminated 60-80% memory waste in LLM inference and became the universal standard for production serving.
Jul 9, 2026
Triton: Block-Level GPU Programming Without CUDA
How OpenAI's Triton compiler enables writing fused GPU kernels through block-level programming, automatic memory coalescing, and tile-based execution, eliminating CUDA boilerplate while matching hand-tuned performance.
Jul 8, 2026
BitNet b1.58: Ternary Weight LLMs That Eliminate Matrix Multiplication
How Microsoft Research's 1.58-bit LLM architecture replaces floating-point matrix multiplications with integer additions, matching full-precision performance while fundamentally changing inference hardware requirements.
Jul 8, 2026
FlashAttention-3: Warp Specialization and the 75% FLOPS Barrier on Hopper GPUs
How FlashAttention-3 exploits H100 warp specialization, asynchronous pipelines, and FP8 quantization to push attention throughput past 75% of peak Tensor Core FLOPS.
Jul 8, 2026
J-Space: Anthropic Discovered a Global Workspace Inside Claude
Anthropic found that Claude has developed a privileged internal workspace, the J-space, where concepts light up silently during reasoning. It mirrors Global Workspace Theory from neuroscience, and it lets researchers read what the model is thinking but not saying.
Jul 8, 2026
Mamba-2: State Space Models and Linear-Time Sequence Modeling Without Attention
How Mamba-2's structured state space duality (SSD) framework unifies SSMs with attention, achieving Transformer-quality language modeling at linear time complexity through hardware-aware block decomposition on modern GPUs.
Jul 8, 2026
Mixture of Depths: Dynamic Token Routing for 50% Faster Transformer Inference
How Google DeepMind's Mixture of Depths achieves equivalent language model quality at a fraction of the FLOPs by learning which tokens can skip entire transformer layers, yielding up to 50% faster inference with a static computation graph.
Jul 8, 2026
RadixAttention: Prefix-Aware KV Cache Reuse for 10x LLM Serving Throughput
How SGLang's RadixAttention uses a radix tree to automatically detect and reuse shared prefixes in KV caches, eliminating redundant computation and achieving up to 10x throughput gains for multi-turn LLM workloads.
Jul 8, 2026
Ring Attention: Distributing Million-Token Contexts Across Devices
How Ring Attention eliminates the memory wall for long-context transformers by overlapping blockwise attention computation with KV-cache communication in a ring topology, enabling near-linear context scaling across devices.
Jul 6, 2026
EAGLE-3: Why the Best Draft Models Stopped Predicting Features
EAGLE-3 gets up to 6.5x decoding speedup by abandoning the feature-prediction objective that defined its predecessors. The interesting part is why feature prediction became the bottleneck, and how a trick called training-time test fixes the train/inference mismatch it leaves behind.
Jul 6, 2026
Native Sparse Attention: Why Trainable Sparsity Beats Post-Hoc Pruning
DeepSeek's NSA makes attention sparsity a first-class citizen of pretraining instead of an inference-time hack, and pairs it with a kernel design that actually turns theoretical FLOP savings into wall-clock speedups. A close read of the architecture and why most sparse attention schemes before it failed to deliver.
Jul 6, 2026
RaBitQ: 32x Vector Compression With an Error Bound You Can Actually Prove
Product Quantization has powered billion-scale vector search for 15 years, but it can fail badly on real datasets and offers no theoretical guarantees. RaBitQ (SIGMOD 2024) compresses vectors to one bit per dimension, estimates distances with a popcount, and comes with a provable O(1/sqrt(D)) error bound. Here is how a random rotation makes that possible.
Jul 5, 2026
How I Built This Portfolio (and Its Blogging Engine) in Angular 22
A technical walkthrough of rebuilding my portfolio from Angular 9 to Angular 22 with spartan-ng, Tailwind v4, SSG prerendering, a WOW animation system, and a zero-backend blogging engine powered by markdown files.
Jul 5, 2026
Two Workloads in a Trench Coat: Prefill/Decode Disaggregation in LLM Serving
Prefill and decode have opposite hardware profiles, and serving them on the same GPUs wastes both. A practical tour of DistServe and Mooncake, the two papers behind the biggest architecture shift in LLM inference.
Jul 9, 2025
Test-Time Compute Scaling: Search, Verification, and the Reasoning Frontier
Why spending more compute at inference time, through process reward models, beam search over chains-of-thought, and Monte Carlo tree search, can outperform simply scaling model parameters. The architectural patterns behind o1, DeepSeek-R1, and compute-optimal reasoning.