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.
55 posts · ~406 min of reading
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.