Independent research and consulting in HPC performance optimization, GPU computing, and AI infrastructure. We profile, analyze, and optimize at the intersection of hardware and software.
Multigrid is the rare algorithm that is provably optimal: it solves a discretized elliptic PDE in a number of cycles independent of the mesh size, for O(N) total work. The idea behind it is a single observation — that error has a spectrum, that cheap iterative solvers only kill the oscillatory part of it, and that a smooth error becomes an oscillatory one if you look at it on a coarser grid. This post walks that idea from the spectral picture to the V-cycle, then makes an argument: multigrid is a *better* fit for a 2026 accelerator than for the machine it was invented on, because the scarce resource has flipped from arithmetic to bytes-moved and synchronizations-taken. Cheap FLOPs rehabilitate exactly the components a FLOP-counting cost model threw out — polynomial smoothers, matrix-free high-order operator evaluation, redundant coarse solves, parallel-in-time — and the same repricing is bringing back a whole family of matrix-theory discretizations. Stated precisely: cheap FLOPs did not make multigrid compute-bound; they moved its optimal design point to where the version worth running is. Includes a complexity-ladder comparison — O(N) multigrid against the O(N log N), O(N^1.5), and O(N^2) alternatives, why each exponent is what it is, and how scaling up (minimize the exponent) diverges from scaling out (minimize communication), mapped onto the kinds of parallelism each kernel wants. To test that, I built a roofline model of matrix-free high-order operator evaluation, then rented an A100 and measured it. The model predicted a sweet spot at polynomial degree 5; the hardware agrees, at 107 ps/DOF — and the measurement shows the model was right by luck. Its flop count was 3x low, arithmetic intensity exceeds the ridge at every degree including p=1, and at the optimum the kernel runs at 5% of DRAM peak and 85% of shared-memory throughput. The DRAM roofline never binds; the sweet spot is set on-chip. Closes on what the idea teaches a loss function, where multigrid has genuinely landed three times over — in the architecture, in the operator, and in the training schedule.
A physics-informed neural network turns a PDE into a loss function — a genuinely elegant idea, and one that is easy to oversell. This post introduces a new interactive explainer from the plasma-PINN campaign that walks the whole arc in diagrams and equations: the canonical formulation and the network wiring, the physics-informed loss, the promised niche of inverse problems, and then the reckoning — a mesh beats the PINN by four to five orders of magnitude on the forward solve, and a convergence proof (Doumèche, Biau & Boyer) shows that an unregularized PINN can drive its training loss to zero while its true error runs to infinity. The through-line is a plain is/is-not ledger: a PINN is a convenient, mesh-free way to *pose* a physics problem — not, by itself, a fast or trustworthy way to *solve* one. What makes it trustworthy is the regularizer, and the regularizer is portable to cheaper machinery.
The optimal way to split N devices into data × tensor × pipeline parallelism is not a constant — it depends on the workload, and the workload changes. As sequence length grows, activation memory grows with it and eventually stops fitting, forcing more tensor/pipeline sharding and leaving less room for data parallelism. So the best config morphs: (64,1,1) at short sequences → (8,8,1) → (4,8,2) at long ones. A controller that re-picks each phase stays on the throughput ceiling; a static config either OOMs later or wastes 16× of its data parallelism early. This is the endgame of the whole track: re-choosing the split mid-run is closed-loop control over the network's own topology.
The previous posts built the traffic matrix M_ij in simulation. This one closes the loop on real hardware: recover a measured M_ij from NCCL Inspector's JSONL, then use its structure to choose between a cheap rail-optimized fabric and a full fat-tree. The catch that makes it interesting — an Inspector record is per-communicator, not per-peer, so reconstructing rank-by-rank traffic means joining Inspector's volumes with your parallelism layout and overlaying the collective's algorithm. Once you have M_ij, the fabric decision falls out of its structure, not its volume: a dense TP+DP job is all on-rail (rail-optimized wins at 1/10th the cost), while an MoE all-to-all of identical volume saturates that same fabric 7.5x over. Runnable end-to-end on a laptop, drop-in for a real NCCL_INSPECTOR_DUMP_DIR.
A synchronous ring all-reduce has a brutal property: if one rank dies, the whole collective hangs — every survivor blocks at the barrier forever. So at scale, where failure is the steady state, membership becomes a live variable and the ring is re-formed on every failure and join. That makes the heartbeat timeout a genuine two-sided tradeoff: too long and every real failure hangs the cluster for ~timeout before it's even detected (detection latency IS stall time); too short and a transient GC pause on a healthy node is mistaken for death, triggering an unnecessary eject-and-rebuild. The total is a U-shaped curve with a clear optimum, and goodput collapses as failures get more frequent — the pressure behind asynchronous training and cheaper reconfiguration.
Two experiments on the same NVIDIA L4, from two different worlds — a JAX @jit on a machine-learning kernel, and an OpenACC data region on a scientific-computing stencil — each turned a one-line change into an order-of-magnitude speedup. Neither did less arithmetic. Both moved less memory. JAX's jit gave 33.9x by letting XLA fuse an elementwise chain into a single pass over HBM instead of materializing every intermediate; OpenACC's `#pragma acc data` gave ~40x by keeping a stencil's arrays resident on the GPU across 2000 iterations instead of shuttling them over PCIe every step. The lesson is the compute-layer lens made concrete: modern accelerators are so fast at math that the tuning is almost never about compute — it's about where the data lives and how many times it crosses a wire.
An all-reduce is a contract — 'every rank ends with the elementwise sum of every rank's input' — not an algorithm. Three different algorithms satisfy that contract (naive gather/broadcast, bandwidth-optimal ring, latency-optimal recursive-doubling), and which one is cheapest morphs with the world size N and message size M. That crossover is exactly why tensor-parallel and data-parallel all-reduce, in the same Megatron job, end up using different algorithms. Built from scratch over real OS processes, measured in bytes-on-wire and message rounds, with the (N, M) crossover surface plotted.
At scale a neural network and the cluster running it are the same object — and that object's communication structure is dynamic. The wiring changes per-step (MoE routing decides who talks to whom, per token), per-phase (training all-reduce vs decode KV movement), and per-event (a node dies, you scale out). This is the opening post of a new track that studies neural nets through the lens of distributed systems: every primitive a large model 'rediscovers' — collective routing, consistency models, load balancing, failure detection, adaptive reconfiguration — is a decades-old idea from distributed systems, renamed. CPU-first, runnable, with Megatron-LM as the production north star.
Training frameworks and serving engines are organized around different things, and for serving the organizing principle has a name: the KV cache. This post starts from the tool map — who shards a training step, who packs sequences into HBM, who routes across a fleet — then goes deep on the one structure that governs serving: what the KV cache is, how it differs from weights, activations, and compute, why decode is bandwidth-bound because of it, and what every serving engine (PagedAttention, RadixAttention, quantized KV, GQA/MQA, offload, disaggregation) is ultimately doing about it. The one-line takeaway: weights decide whether the model fits; the KV cache decides how many requests you run at once and how fast each decodes.
The question 'how fast is this LLM system?' has three different answers, and conflating them is the most common mistake teams make. This is a field guide: the three layers of benchmark (system performance, model capability, agentic system performance), how MLPerf brings rigor to the first, and how the unit of work is shifting from a single prompt to a 200-turn agent trajectory — the change Artificial Analysis's AA-AgentPerf is built to measure, with 'agents per megawatt' as its headline metric. The throughline: agentic performance is our same compute-layer story, restacked.
Synchronous SGD keeps the parameter linearizable: every worker sees the exact global average every step. That barrier is also the straggler tax — the whole job runs at the speed of its slowest worker. Distributed datastores hit this wall years ago and answered with weaker consistency; SGD rediscovers the identical knob. This post builds a CPU-only event simulator of N workers with one straggler and dials consistency from sync (BSP) through bounded-staleness (SSP) to fully async (ASP), measuring the real tradeoff: weaker consistency buys throughput and costs per-update progress, and the staleness bound s clips the tail, not the mean.
LLM inference comes wrapped in an alphabet soup — DP, TP, PP, EP, SP, CP, KV cache, prefill/decode, TTFT, TPOT, MoE. This is a practitioner's glossary that organizes the acronyms by the problem each one solves rather than alphabetically, because they're not independent terms — they're all answers to one question: what do you do when a model won't fit, or won't serve fast enough, on a single GPU? Inspired by NeuReality's 'guide to the salad of acronyms,' extended with the terms my own measurements keep running into.
In a mixture-of-experts layer a learned router decides, per token and per batch, which expert — hence which device — each token is sent to. So the all-to-all that dispatches tokens has a traffic matrix that is a function of the input, recomputed every batch. This is the sharpest example of the network morphing during activity: you cannot draw a static wiring diagram. A CPU-only simulator shows skew driving expert imbalance, token drops, and the all-to-all bottleneck up in lockstep — and shows the busiest device wandering batch to batch, which is exactly why production MoEs lean on a capacity factor and an auxiliary load-balancing loss.
Backpropagation is from the 1980s, so why did deep learning erupt only in the last decade? Not because the math changed — because three things converged: self-supervision lifted the data ceiling from 'what we can label' to 'what exists'; the GPU turned out to be a throughput machine shaped exactly like the workload (thousands of simple in-order ALUs that hide memory latency behind arithmetic, fed by high-bandwidth memory); and the Transformer is GPU-shaped by design. LLMs are just the point where that loop was scaled until the systems became the whole game. A foundational explainer, with the hardware intuition made precise.
I ran the comparison everyone asks for — vLLM vs SGLang, same GPU, same model, same workload, same harness — expecting a throughput winner. Instead, on vanilla single-shot generation the two are neck-and-neck: within single-digit percent across the whole concurrency frontier. The real differences live elsewhere — hardware portability (vLLM ran on a T4; SGLang wouldn't) and workload shape (where SGLang's RadixAttention should pay off, though a short shared prompt gave it little to bite on here — see the in-post correction). A lesson in not picking an inference engine by a leaderboard number.
Most LLM code is Python, and most of the time a traceback is enough. The bugs that stop a training run or a serving deployment cold live below the framework — in the native CUDA, C++, and collective layers where a Python stack goes dark. This post walks one deliberately-broken workload down the compute ladder — gdb on a CPU segfault, compute-sanitizer and cuda-gdb on a GPU out-of-bounds — with real transcripts captured on a T4, then names the tools for the rungs above it: rocgdb on AMD, gdb4hpc across MPI ranks, and Linaro DDT at scale.
Our last post filled a single GPU with batching — scaling *up*. The next question is scaling *out*: when one GPU's throughput ceiling is reached, you add replicas, and Kubernetes is the orchestrator that does it automatically. We build the whole stack on a real T4 — k3s, the NVIDIA device plugin, time-slicing, vLLM, and a HorizontalPodAutoscaler — and watch it scale 1→2 replicas under load and back down. The honest finding: on a single physical T4, time-sliced replicas share the silicon, so the autoscaling control loop works perfectly while aggregate throughput stays GPU-bound. Scale-out is not scale-up. Plus three production bites — a BF16 cast, host RAM exhausting before VRAM, and a 4-vCPU control plane starving — that separate a slide from a deployment.
Our memory-bandwidth study found that high-bandwidth memory only pays off with enough parallelism to fill it. LLM inference is that principle at production scale: token-by-token decode is memory-bandwidth-bound, and batching is the parallelism that unlocks the GPU. We measure it directly — vLLM serving Qwen2.5-1.5B on a single T4 — and watch throughput climb 27x from concurrency 1 to 64 while latency degrades on a predictable frontier. This kicks off a new line of work on LLM systems.
The two gem5 studies modeled one engine and its memory system. Modern silicon is no longer one engine — it is a CPU, GPU, and NPU sharing a fabric, and the hard problem has moved to the boundaries between them. This post argues that heterogeneous-SoC performance is, formally, a models-of-computation interface problem — one Axel Jantsch wrote down in 2003 — and that cycle-level pre-silicon simulation is the instrument for reasoning about it before the chip exists.
A follow-up to our cache-size study: we measure memory bandwidth instead of latency on a simulated ARM system, and find that the fastest DRAM is not the fastest in practice. The deciding factor is parallelism — the same DDR3 vs DDR5 comparison flips its verdict depending on whether one core or a flood of requests is driving memory. This is the GPU/HBM rationale, measured from first principles.
A reproducible system-level investigation: we sweep L2 cache size on a simulated ARM core and find that the cache that maximizes performance is not the cache that minimizes energy. Using gem5 for cycle-level timing and McPAT for power, we locate the energy-optimal sweet spot and explain the U-shaped energy curve that leakage creates.
A structured practice repository for running multiple Claude agents in parallel — Zed threads and the async Python SDK — that grew into something bigger: a fleet of independent Claude Code sessions, each in its own terminal and git working directory, coordinated by a written ownership map, a resident 'main office' coordinator agent, and an append-only check-in ledger. This website is what that multi-agent system produced.
Documenting the profiling methodology and toolchain we'll use across our AI framework performance investigations. Covers Nsight Systems, Nsight Compute, PyTorch profiler integration, and custom instrumentation approaches.
III. Focus Areas
Pre-Silicon & Whole-System Co-Design
Exploring novel concepts pre-silicon — co-designing algorithm, runtime, and architecture before the hardware exists. A holistic, whole-system view, where the performance win comes from the full stack working together rather than any single layer.
AI Framework Performance
Profiling and optimization of PyTorch, NeMo, Megatron-Core, and inference runtimes (vLLM, TRT-LLM). End-to-end training and serving pipeline analysis.
GPU Architecture Analysis
Deep performance characterization on NVIDIA and AMD GPU architectures. Memory hierarchy analysis, occupancy tuning, kernel optimization using Nsight and ROCm profiling tools.
HPC & Distributed Systems
MPI, NCCL/RCCL communication optimization. Multi-node scaling analysis. Performance engineering for exascale-class systems.
Custom Instrumentation
Beyond vendor tools — custom profiling, roofline modeling, memory bandwidth characterization, and performance modeling for novel workloads.
Systems & Concurrency
OS and runtime effects that gate accelerators — memory allocation, NUMA placement, page faults — plus thread contention and data-race correctness across precisions and GPU architectures.