Investigations

Performance and correctness studies in GPU computing, HPC, and AI systems — profiling, benchmarking, and instrumentation, with reproducible code and data.

Know, Do, Decide: Where an Agent Lives on the Computing Landscape

Ask a plain language model why pickles are wiggly and it answers in about a second, entirely on the GPU: one prefill, ten decode steps, done. Ask an agent the same question and you get three GPU visits separated by long stretches of CPU work and waiting — roughly three times the tokens and four times the wall-clock to reach the same sentence, because the model already knew the answer. That comparison is the clearest way to place an agent on the computing landscape: it is not a new hardware tier at all. It is a control loop that runs on the CPU, rents the GPU by the lap, keeps its memory as re-fed text rather than in any silicon, and spends most of its life waiting on the world. Which tells you exactly when to use one, and what to optimize when you do.

Metering Goodput: The Number Your Dashboards Don't Have

Goodput — throughput that actually meets your SLA — is the metric that decides whether users are served, and it lives everywhere except where operators look: benchmark harnesses have it (NVIDIA AIPerf's --goodput flags, vLLM's bench), papers formalized it (DistServe), Google ships a training-side library — but Prometheus, Grafana, and DCGM dashboards have no goodput concept at all. Production fleets alert on raw tok/s and error rate, the two numbers our own measurements show staying green while service collapses. So we built the missing piece: a ~100-line always-on goodput meter — cache-busted synthetic probes, rolling SLA conformance, error-budget burn, Prometheus textfile out — and in its first 60 seconds it caught a cold-load and a load spike that raw throughput never saw. The gap between benchmark-time and always-on metering is small enough to close in an afternoon; what's scarce is deciding to state an SLA and alert on conformance.

Will the Models Be Shelved? Copyright, Corpora, and the Price of Training Data

The largest copyright settlement in U.S. history — Anthropic's $1.5B, finally approved July 2026 — ordered the destruction of a pirated dataset, not of any model. Meta won fair use on the record presented. The New York Times case, which explicitly asks for model destruction, heads to summary judgment. So which is it: do models get shelved, or do companies pay and keep serving? The emerging pattern is precise about where liability attaches — acquisition and outputs, not the weights — which so far makes copyright a cost, not an existential threat. But shelving is not hypothetical (a legal-AI startup was litigated out of existence; the FTC has ordered models destroyed four times), the statutory-damages math has a tail that changes everything, and the structural fix is the one the industry is slowest to adopt: corpora with provenance.

Claimed, Packed, Working: What Using the Hardware Actually Means

Three meters all sound like 'how hard is the GPU working' — and the two everyone watches can read 100% while the one that matters reads 7%. nvidia-smi utilization measures whether the chip is claimed. Occupancy measures whether its scheduler is packed. Neither measures whether it is working: an H100 running a competent plain-CUDA matmul is claimed, packed, and ~93% idle, because the tensor cores only exist for kernels that issue their instructions. This post climbs the whole ladder — from dark silicon through MFU/MBU, goodput, and $/solved-task — and lands on the only metric a buyer actually feels: solutions per dollar per hour. Every rung is a place the fleet quietly lies to you, and every rung has been measured somewhere in this practice's work.

The Agent Loop, Pedantically: What an AI Agent Can and Cannot Do

An AI agent is not a new kind of model. It is an ordinary next-token predictor wrapped in a plain while-loop by a harness, and every mysterious-sounding agent behavior — planning, deciding, 'calling' a tool — is just tokens emitted into that loop. Here are the seven pedantic steps of one iteration, stated with no hand-waving, and the honest capability map that falls out of them: an agent can do whatever the loop lets it observe and the tools let it touch — and it cannot execute anything itself, remember beyond the re-fed context, learn mid-session, or verify itself without an external check. Most agent-engineering folklore (context budgets, verifier-gated routing, permission layers, turn caps) stops being folklore once you see which step it defends.

Do Agents Run on CPU or GPU? Both — and Mostly Neither

A simple question with a three-part answer that falls exactly along the agent loop's seven steps. The harness — context assembly, tool execution, the loop itself — is ordinary CPU work. The model's forward passes live on the GPU, though decode-dominated agent workloads are memory-bandwidth-bound, which puts them in the regime where a matrix-enhanced CPU with fast memory is genuinely competitive. And the wrinkle that matters most in practice: an agent's wall-clock is often dominated by neither processor, because it is waiting on the world — a test suite, an API, a build. That makes serving agents a scheduling problem as much as a hardware one, and it is why a bake-off can find an A100 barely stressed while the agent on top of it works flat out.

Trailing the Frontier: How the GPU Cloud Actually Makes Money

About 90% of leading GPU revenue now comes from AI, and the newest chips are designed around it — enormous low-precision throughput at the direct expense of the double-precision math that predictive science depends on. Vendors are cutting native FP64 and pushing emulation instead. Meanwhile the GPU cloud makes its recurring margin not on the frontier but on the trailing edge: inference on depreciating hardware. This is the accuracy-versus-cost frontier as a business model — and why the scarce, valuable thing is knowing what accuracy you're actually getting, at what cost.

The General Absorbs the Specific: Why the CPU+GPU Tandem Survives the 'Do We Still Need GPUs?' Question

In June 2026 a GPU-less Chinese supercomputer, LineShine, took TOP500 #1, and three of HPC's most credible names — Dongarra, Matsuoka, Hoefler — published a paper asking whether the discrete GPU is still necessary now that CPUs have absorbed wide vectors, matrix engines, on-package HBM, and every precision from FP64 to FP4. The honest reading of that paper is not 'kill the GPU.' It is a convergence thesis, and convergence is the opposite of elimination. Here is the breakdown, and the argument for the don't-throw-anything-away posture: route each kernel to the substrate that fits its roofline, and treat CPU+GPU as one heterogeneous machine — multigrid in hardware.

Why Benchmark Scores Lie (and What to Demand Instead)

A frontier model scores 85–90% on the standardized test and then delivers consistent output less than a quarter of the time on the same task in production. That gap is not noise — it has two nameable causes that two independent 2024–2026 sources converge on: the benchmark is often badly built (Stanford's BetterBench graded 24 popular benchmarks and found only 3 ship a runnable replication script), and even a well-built benchmark's score is inflated by contamination and weak construct validity (the industry's test-vs-reality reliability gap). What both are really telling you, and the checklist to demand before you trust a number.

Beyond Tokens: The Model Input/Output Zoo

The text-LLM loop — symbols in, one random draw per token, symbols out — is a single point in a much larger design space. Drop the constraint that input and output are both human language and you get every other model family: continuous signals in (audio, vision, coordinates), structured inputs (graphs, tables), continuous outputs (images, fields, actions, embeddings), and mechanisms that replace the autoregressive token loop entirely. A field guide, organized by where symbols and numbers trade sides.

It's Numbers All the Way Through: From Prompt to Token

Your prompt starts as text and ends as text, but in between it is nothing but numbers. The tokenizer turns words into integers; a lookup turns integers into vectors; and from there until the very last step the model is pure linear algebra. There are exactly two places where symbols and numbers trade sides, and exactly one place where anything random happens. A visual walk through the whole pipeline, with a detailed look inside the tokenizer's byte-pair merge loop and how the same word tokenizes differently across two vocabularies.

The Tokenizer Tax: Theory, Bottlenecks, and the Loss It Shapes

Tokenization looks like plumbing, but it is a compression scheme that quietly decides how much a prompt costs, what units the loss is measured in, and an accuracy ceiling on anything below the token boundary. The theory (BPE, WordPiece, Unigram), the known bottlenecks (fertility, character blindness, glitch tokens, head cost), the speed-up strategies (fast tokenizers, bigger vocabs, tokenizer-free models), and the part almost no one explains: the connection to per-token cross-entropy, bits-per-byte, and accuracy.

The Denoising Engine: Diffusion Models, from Basics to Scale

A diffusion model turns generation into iterative denoising — and every beginner primer ends on the same limitation: sampling is slow. This review starts there, from the fundamentals (forward noising, reverse denoising, the noise-prediction objective), and then does the part the primers skip: a serious walk through how the field drove sampling from ~1000 steps to one. Better ODE solvers, latent-space diffusion, trajectory distillation and consistency models, cross-step feature caching, quantization — then the scale-up story, where the U-Net becomes a Transformer, a single image or video is parallelized across GPUs, and diffusion quietly merges with the LLM serving stack.

GPU-BLAST on T4: A 1% Roofline and What It Means

GPU-BLAST v1.1 on a Tesla T4 achieves 2.13× over one CPU thread — but four CPU threads win outright at 3.2×. The roofline explains why: arithmetic intensity is 12 ops/byte (memory-bound regime), but the kernel lands at 1% of the bandwidth ceiling. The gap isn't bandwidth — it's DRAM latency. BLAST's scoring matrix lookup is one of the worst-case access patterns for a GPU: two data-dependent indices, zero spatial locality, and only 2 warps per block to hide 400-cycle stalls. The fix is four lines of shared memory code.

Measuring Training Goodput: What a Killed Rank Actually Costs

Goodput at scale is an HPC-systems problem, not an ML-modeling one: useful-work-per-dollar is set by failures, restarts, stragglers, and collective tails. This post measures one slice honestly on 8×A100 — what a single killed rank costs. The headline is counterintuitive: killing a rank does not crash the job, it strands it. Detection is fast (~9 s) but the seven surviving ranks hang in the collective for ~270 s, burning GPU-hours, before teardown; that avoidable hang dominates the ~6-minute per-failure cost, not the checkpoint gap. A separate straggler experiment shows one GPU at half clock taxes all eight ranks by 70% — a synchronous step runs at the speed of its slowest rank. Two honest ops findings ride along: the nsys profiler moves the very metric it measures by ~24% while active, and a naive checkpoint-restart took three config fixes to work. Closes on the fleet arithmetic a CSP goodput engineer owns.

Goodput: The Only Number That Pays the Bill

MFU tells you how hard the machine is working. Goodput tells you how much of that work you get to keep. The gap between them is failures, rollbacks, stragglers, and jitter — and at fleet scale that gap is where the money goes. A run at 45% MFU that loses an hour to a dead rank every six hours is not a 45% run. This post defines goodput, walks the arithmetic of rollback (why shrinking MTBF eats you alive as N grows), and argues the uncomfortable part: goodput is an HPC systems problem — fault tolerance, collective resilience, straggler and noise mitigation — not an ML-modeling problem. The resilience layer under a GPU fleet is systems, not PyTorch.

Dataflow vs Quantum: How I Understand Two Post-von-Neumann Bets That Aren't Competing

"New computing" gets sold as one undifferentiated bucket — quantum, dataflow, neuromorphic, all filed under 'the future of the chip' — and untangling that bucket is how I came to actually understand what each of these machines is for. This is my working model of how dataflow and quantum relate, built by connecting both back to things I already reason about: the roofline, Little's Law, and — because I came up through plasma physics — the fusion problem. Dataflow and quantum are both answers to 'how is the machine model changing,' but they change different things and win in different ways: dataflow changes the execution trigger (control-flow to data-availability, temporal to spatial) for a constant-factor throughput/energy win that raises the classical roofline; quantum changes the representation of state itself (bits to qubits, deterministic to sampled) for — in a narrow slice of problems — a complexity-class change no amount of classical silicon can match. I use the July 2026 ORNL/Cleveland Clinic/IBM fusion-materials result (FLiBe tritium-breeding chemistry on IBM Heron QPUs) as the concrete thing to reason about, and treat it with the same skepticism this lane points at vendor benchmarks. Where I've landed: they're complementary, not competing — every hybrid quantum algorithm and every error-correction loop needs a low-latency classical throughput partner, which is exactly the shape a dataflow fabric is built to eat. The framing that finally made it click for me isn't 'dataflow vs quantum'; it's 'dataflow as the classical half of a quantum-classical machine.'

The Dataflow Landscape: How I Learned to Place NextSilicon Among Its Neighbors

When I first tried to understand NextSilicon's Maverick-2 I made the beginner's mistake of studying it alone, as if it were a singular invention. It isn't — it's one point in a crowded, decades-deep space of spatial and dataflow architectures, and I couldn't judge what was special about it until I had a map of the neighborhood. This is the map I built for myself, organized around the one question that finally made the whole field sort itself: when is it decided which operation runs on which cell? That single axis — from dynamic and runtime-reconfigured (NextSilicon) through compile-time reconfigurable (SambaNova, FPGAs) to static and deterministic (Groq) to fixed-function (the TPU's systolic array) — turned a confusing pile of vendor names into a spectrum I could reason about. I walk the commercial field, the academic ancestors NextSilicon productizes (Plasticine, TRIPS/EDGE, WaveScalar), and the systolic array that quietly won production. Three things I understand now that I didn't before: more scheduling dynamism buys generality but costs overhead; NextSilicon sits in a nearly empty cell (the only one that's both runtime-reconfigurable and runs unmodified general HPC code); and the most successful dataflow chip in the world is also the least reconfigurable — which kept me honest about the whole thesis. All capability and performance claims are vendor positioning, flagged as such.

Inside the Fabric: How I Came to Understand NextSilicon's Dataflow Machine

This is my attempt to understand a genuinely new kind of processor by connecting each piece of it to something I already knew. A reconfigurable dataflow chip makes an audacious promise — compile your existing HPC code once, and a grid of ALUs rewires itself at runtime to run your hot loops from live telemetry — and the only way I could make it stop sounding like magic was to keep asking 'what does this remind me of?' The answer, four times over, was a machine I'd met before: the way it runs hundreds of threads per core is a barrel processor; the way it feeds the grid from memory is decoupled access/execute; the way it survives branches is if-conversion; the way it rewires itself is a tracing JIT. Once I had those handholds, the whole thing resolved into a single idea — it's a Little's-Law latency-tolerance machine whose one genuinely new trick is the telemetry-driven runtime remap. To keep myself honest I pulled three of NextSilicon's patents, which caught me being wrong twice (the famous patent is about multithreading, not telemetry; the front-end is not Tomasulo) and confirmed one guess I'd flagged as a guess (the telemetry counters are claimed, verbatim, in the fabric patent). I've left the corrections in, because that's how the understanding actually got built. Every vendor performance number is a claim; every undisclosed constant is flagged; and this is my current understanding, not the last word.

Multigrid: The O(N) Idea, and What It Teaches a Loss Function

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.

What a PINN Is — and Is Not: A Visual Field Guide

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.

Adaptive Parallelism: Closed-Loop Control Over the Topology

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.

Seeing the Traffic Matrix: From NCCL Inspector to a Fabric Decision

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.

Elastic Membership: When Failure Is the Steady State

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.

One Line, an Order of Magnitude: Why the Speedups Live in Data Movement

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.

All-Reduce, Three Ways — and Why the Cheapest One Keeps Changing

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.

Your Neural Network Is a Distributed System That Won't Hold Still

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.

What the KV Cache Actually Is — and Why It's the Binding Constraint

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.

From MLPerf to Agents per Megawatt: How We Benchmark LLM Systems Now

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.

Consistency Models for SGD: What a Straggler Teaches You About Eventual Consistency

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.

The Salad of Acronyms: A Working Glossary of LLM Inference Parallelism

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.

MoE Routing: The Network Topology That Won't Hold Still

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.

Why Deep Learning Works Now: The Same Math, the Right Machine

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.

vLLM vs SGLang: An Honest Bake-Off (and Why the Numbers Aren't the Point)

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.

Below the Framework: Debugging LLM Workloads from a CPU Segfault to a GPU Out-of-Bounds

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.

Scale-Out Meets the Silicon Ceiling: Kubernetes LLM Autoscaling on a T4

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.

Batching Is the Parallelism: Measuring LLM Inference on a Commodity GPU

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.

Modeling the Heterogeneous SoC: Models of Computation Meets Pre-Silicon Simulation

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.

Fast Memory Isn't Fast Alone: Bandwidth, Parallelism, and the HBM Question

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.

Bigger Isn't Better: Finding the Energy-Optimal Cache with gem5 and McPAT

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.

Parallel AI Agents: From a Practice Repo to the System That Built This Site

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.

Establishing a Baseline: AI Framework Profiling Methodology

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.