The Complete Mental Model, Learning Path, and Projects
Part 19 of 19
This is the last post of Deep Learning Inference Engineering — nineteen parts, ten lessons, one discipline. Lesson 1 opened by asking what it even means to run a trained model, counted in FLOPs and bytes moved; the chapters since have walked that question down through quantization, kernel engineering, memory planning, compilers, and real hardware. This part closes the loop: it collapses everything into one mental model, turns that model into a checklist you can run against any network you're handed, and hands you a learning path and six projects for turning the reading into skill you actually own.
1. Nineteen Parts, One Discipline
Before building anything new, compress what's already here.
Lesson 1 established that FLOPs alone don't predict latency — arithmetic intensity does, and the Roofline model turns that into a line you can plot a kernel against to know, before writing a single instruction, whether you are fighting the ALUs or the memory bus.
Lesson 2 took that same physics and applied it at graph scale: tensor layout as a choice about which logical neighbors become physical neighbors, tiling as the mechanism that keeps a working set inside a memory tier long enough to be reused, and fusion and constant folding as ways to delete work that never needed to happen at runtime.
Lesson 3 went underneath the graph into the numbers themselves — the affine quantization map, why weights are usually symmetric and per-channel while activations are usually asymmetric and per-tensor, what actually happens inside an int8 multiply-accumulate, and how calibration and quantization-aware training decide the range that whole scheme depends on.
Lesson 4 went underneath the graph a second time, into the loop nests — nine largely independent dimensions of kernel optimization, from loop ordering and blocking through register tiling and SIMD, worked through a real GEMM and a real im2col convolution.
Lesson 5 answered the two questions a graph and a set of fast kernels still leave open: where does every tensor actually live, and who decides when and on which processor each operator runs. Static memory planning turned out to be interval scheduling and graph coloring in a trench coat; the runtime turned out to own scheduling, device placement, and overlapping DMA with compute.
Lesson 6 took those runtime mechanics and pointed them at a specific, consequential workload: serving. Latency and throughput pull against each other through batch size, and transformer decoding turned out to have its own particular shape — the KV cache, and the split between a compute-bound prefill and a memory-bound decode.
Lesson 7 zoomed out to the compiler that makes all of the above automatic instead of hand-written per model: why no single IR can be simultaneously close to a framework's semantics and close to hardware, MLIR's dialect stack as the concrete answer to that tension, and layout selection and kernel selection as decisions that only make sense at specific points in that stack.
Lesson 8 asked a different kind of question — not how to run a given model faster, but how to make the model itself smaller and cheaper without asking permission from the runtime: pruning, distillation, low-rank factorization, and structured sparsity.
Lesson 9 supplied the discipline that makes every other lesson's reasoning trustworthy instead of folklore: profiling and benchmarking, replacing a guess about where the time goes with a measurement of where the time actually goes.
Lesson 10 took all of it to the edge, where the constraints stop being theoretical: Amdahl's Law as a spending plan for a fixed latency budget, energy per inference as a currency that sometimes matters more than milliseconds, and microcontroller-class inference where a single unplanned allocation is a hard fault with no operating system standing by to catch it.
Ten lessons, and not one of them introduced a genuinely new kind of problem. Every single chapter was an answer to one of five questions — how much compute, how much data movement, how much parallelism, what precision, and when and where should it run. That's not a retroactive theme bolted on for a tidy ending; it's the actual shape of the field, and Section 4 below makes it precise. But first, two sections that this reader in particular should sit with, because they are not generic connective tissue — they are the fastest path from where you already are to genuine expertise here.
2. How This Connects to a C++ Background — Made Concrete
It is tempting to wave at "inference engineering rewards systems programmers" and move on. That's true, but vague, and vague is useless. Here is the actual mapping, concept for concept, to the specific things a C/C++ systems background already trained into you.
| C++ / systems concept | Inference-engineering concept | Where in this series |
|---|---|---|
Manual memory management — malloc/free, arenas, no garbage collector to bail you out | Static memory planning: computing every tensor's buffer offset once, ahead of time, so the runtime never calls an allocator mid-inference | Lesson 5, Part 1 |
| RAII / ownership discipline — knowing precisely when an object's lifetime ends | Buffer lifetime analysis: a tensor's live range is [birth, death] in graph steps, and getting that interval wrong either corrupts a still-live tensor or wastes memory holding a dead one | Lesson 5, Part 1 |
| Cache-aware data structure layout — array-of-structs vs. struct-of-arrays | Tensor layout selection: NCHW vs. NHWC, chosen so the hardware's native reduction axis lands stride-1 in memory instead of scattered across cache lines | Lesson 2 |
| Pointer arithmetic and stride computation | Address calculation inside a kernel's inner loop — base + (k * N + j) * elementSize is the exact arithmetic a GEMM micro-kernel runs on every iteration | Lesson 4 |
| Loop-nest tuning you'd do by hand in a hot C++ loop — interchange, blocking, unrolling | Kernel blocking/tiling: sizing a sub-block of the computation to survive in a chosen memory tier across every reuse it's going to get | Lesson 2, Lesson 4 |
Placement new, object pools, custom allocators | Arena allocators and fixed tensor arenas — ONNX Runtime's arena allocator and TFLite Micro's head/tail-split arena are the same idea under different names | Lesson 5 |
const correctness and the aliasing rules a compiler needs to vectorize a loop | A compiler's need to prove no-alias before it will fuse two operators or auto-vectorize a loop — the same proof obligation, one level up | Lesson 4, Lesson 7 |
Two of these rows are worth slowing down on, because they are not analogies of convenience — they are the identical mathematical object wearing a different name.
Manual memory management and static memory planning are the same discipline solving the same problem at a different granularity.
When you hand-write a C++ arena allocator for a hot path, you are doing exactly what Lesson 5 showed a memory planner does automatically: computing lifetimes ahead of time, packing non-overlapping objects into shared storage, and refusing to pay for a heap allocator's bookkeeping on a path where you already know the shapes.
The one thing a memory planner adds that a hand-rolled arena usually doesn't bother with is treating buffer assignment as a formal combinatorial problem — the tensor interference graph, where two tensors get an edge exactly when their live ranges overlap, and assigning buffer slots without any two conflicting tensors sharing a slot is graph coloring.
If you've ever manually reasoned about which of two objects in a pool "can't be alive at the same time so they can share a slot," you were doing graph coloring by hand without the vocabulary for it.
RAII maps almost too cleanly onto buffer lifetime analysis to be a coincidence.
RAII's entire premise is that an object's lifetime should be a statically knowable, compiler-enforced interval — construction to destruction, no ambiguity, no manual bookkeeping the programmer might forget.
A tensor's live range in a memory plan is the exact same claim applied to a different kind of object: it is knowable before execution even starts, because inference is a fixed, forward-only DAG with known shapes. Training is the counterexample that proves the rule — backpropagation stretches every activation's lifetime from its forward-pass producer all the way to its backward-pass consumer, and destroys the short-lifetime assumption the whole scheme depends on.
The C++ instinct to distrust any resource whose lifetime isn't provably bounded is precisely the instinct that makes static memory planning legible the first time you see it, instead of feeling like a new idea.
3. How This Connects to Compiler Engineering — Made Concrete
If the C++ mapping is useful background, this one is closer to home. IR design, lowering passes, and instruction and kernel selection are not adjacent to an MLIR-based NPU compiler background — they are that job, described from the inference side of the fence instead of the compiler-internals side.
| Compiler-engineering concept | Inference-engineering concept | Where in this series |
|---|---|---|
| IR lowering through progressively lower-abstraction dialects | MLIR's real dialect stack: framework dialect to TOSA/Linalg-on-Tensor to Linalg-on-Memref/Affine to SCF and Vector to the LLVM dialect | Lesson 7 |
| Register allocation by graph coloring (Chaitin-style) | Static memory planning by tensor-interference-graph coloring — literally the same construction, tensor lifetimes standing in for variable lifetimes, a heap standing in for the register file | Lesson 5, Lesson 7 |
| Instruction selection — choosing which real machine instruction implements an IR operation | Kernel selection — choosing im2col-plus-GEMM, direct convolution, or Winograd to implement one conv2d graph node | Lesson 4, Lesson 7 |
| Loop-nest transformations in an affine/polyhedral compiler — tiling, interchange, fusion, skewing | The Affine dialect's tiling and fusion passes, lowering a linalg.generic op's declared iteration space into an actual blocked, reordered loop nest | Lesson 2, Lesson 4, Lesson 7 |
| Constant folding and dead-code elimination | Constant folding and dead-node elimination applied to the computation graph, plus bufferization's dead-tensor elimination one level lower | Lesson 2, Lesson 7 |
| Autotuning / superoptimization search over a schedule space | AutoTVM- and Ansor-style search over tile sizes, unroll factors, and loop order for one (operator, shape, hardware) triple | Lesson 7 |
The instruction-selection-to-kernel-selection mapping deserves the most weight, because it is the single cleanest one-to-one correspondence in the entire series.
A classic backend's instruction selector looks at a DAG of target-independent IR operations and decides which real machine instruction — or short instruction sequence — implements each one, given the specific target's instruction set.
Kernel selection is that same decision, one abstraction level up: given a conv2d node in graph IR, the compiler has to decide whether it becomes a direct seven-nested-loop convolution, an im2col-plus-GEMM formulation that reuses the same register-blocked micro-kernel every fully-connected layer already uses, or a Winograd transform that trades extra additions for fewer multiplications on small kernel windows.
All three are semantically identical — same mathematics, same output — and differ enormously in which one a real vectorizer and register allocator can turn into fast code on a given target. That is exactly what instruction selection is doing when it picks a fused multiply-add over a separate multiply-then-add: same math, different cost on this specific hardware.
Register allocation and memory planning being the same algorithm rather than a similar one is worth restating plainly, because it's easy to read past it as a cute parallel rather than register the actual claim: build an interference graph over live ranges, color it with the fewest colors, spill to slower storage whenever the graph can't be colored within budget.
Chaitin's 1981 algorithm does this over variable live ranges with a fixed physical register file as the color budget. A memory planner does the identical construction over tensor live ranges with a memory budget standing in for the register file, and spills to a larger, slower tier — DRAM instead of SRAM — under the same pressure a register allocator spills to the stack.
If you have ever debugged why a compiler spilled a hot variable it "should have" kept in a register, you already have the intuition for why a memory planner sometimes fails to hit the theoretical minimum-buffer count on an irregular tensor graph: the underlying obstruction is the same one, dressed differently.
The dialect-stack mapping is the second one worth internalizing, because "progressive lowering" is not an ML-specific invention.
It's LLVM's own AST-to-IR-to-SelectionDAG-to-MachineInstr shape, just stretched across a much larger abstraction gap — the distance from "a PyTorch nn.Conv2d call" to "seven nested loops with explicit index arithmetic" is far larger than the distance LLVM IR alone was ever designed to span.
Every dialect transition deletes information its own level didn't need — a graph-IR pass has no use for "this loop's access pattern is stride-1 in the innermost dimension," and a vectorizer has no use for "this was originally a framework op call" — while adding exactly the information the next level requires.
That is the same discipline behind any well-factored compiler pass pipeline: each pass should operate at the abstraction level its analysis actually needs, not reconstruct high-level structure from low-level detail it was never given.
4. The Five-Dimension Mental Model
Here is the claim the previous two sections were building toward: every optimization in this entire series, without exception, is an answer to one of exactly five questions.
Dimension 1 — COMPUTE
How many operations does this actually require?
Dimension 2 — DATA MOVEMENT
How many bytes have to move, and across which memory tier?
Dimension 3 — PARALLELISM
How much of this work can happen at the same time?
Dimension 4 — PRECISION
How many bits does each value actually need to carry?
Dimension 5 — SCHEDULING
When, and on which piece of silicon, does each operation run?Treat this as a lens, not a taxonomy — most real techniques sit on two axes at once, one dominant and one secondary, and naming both is usually the fastest way to understand why a technique works instead of just that it does. The table below runs a representative slice of this series' actual content through the lens, as a demonstration rather than an exhaustive catalog.
| Technique | Lesson | Primary axis | Secondary axis |
|---|---|---|---|
| Roofline model / arithmetic intensity | 1 | Data movement | Compute |
ikj loop reordering for GEMM | 4 | Data movement | — |
| Register-blocked GEMM micro-kernel | 4 | Compute | Data movement |
| NHWC vs. NCHW layout selection | 2 | Data movement | Parallelism |
| Operator fusion | 2 | Data movement | Scheduling |
| Affine-dialect tiling pass | 7 | Data movement | Compute |
| Per-channel int8 quantization | 3 | Precision | Compute |
| Quantization-aware training | 3 | Precision | — |
| Structured pruning / sparsity | 8 | Precision (information density) | Compute |
| Knowledge distillation | 8 | Compute | Precision |
| NEON / AVX SIMD vectorization | 4 | Parallelism | Compute |
| Multi-stream execution, device placement | 5, 9 | Parallelism | Scheduling |
| Batch size selection | 6 | Parallelism | Scheduling |
| KV cache | 6 | Data movement | Scheduling |
| Static memory planning / graph coloring | 5 | Data movement | Scheduling |
| Zero-copy execution, DMA overlap | 5, 10 | Data movement | Scheduling |
| Amdahl's Law latency budgeting | 10 | Scheduling | — |
| Systolic / MAC-array dataflow | 10 | Parallelism | Data movement |
A few of these are worth pulling out because the mapping clarifies something the original lesson stated but didn't need to frame this way.
The Roofline model is fundamentally about data movement — arithmetic intensity is FLOPs per byte, and the whole model exists to answer "is this kernel starved for bytes or starved for ALUs." But it only becomes actionable once you also know the compute side of the ratio, which is why it's listed with compute as secondary rather than absent.
Pruning and sparsity sit oddly under "precision" until you see it the right way: precision, in the fullest sense, isn't only "how many bits per value," it's "how much information does this value actually carry." A pruned-away weight is the limiting case of a weight that needed zero bits, not four or eight.
KV cache is data-movement-primary because its entire reason for existing is avoiding the re-computation of already-computed key and value projections, which is really avoiding re-reading the tokens that produced them. But it is inseparable from scheduling, because the split between a compute-bound prefill and a memory-bound decode is precisely a scheduling fact about when the bottleneck changes shape mid-request.
The reason this five-axis grid is worth building at all, rather than just remembering "there are five things to think about," is that it gives you a diagnostic move you can run on any technique you haven't seen before: name its primary axis, name its secondary axis, and you almost always already know, from the axis alone, what its failure mode will be and what it will trade away.
A parallelism-primary technique will hit an Amdahl's Law ceiling. A data-movement-primary technique will stop paying off once you're compute-bound. A precision-primary technique will eventually cost accuracy.
That single move — read a technique's axis, predict its ceiling — is most of what separates recognizing a familiar-looking optimization from actually reasoning about a new one.
5. The Questions to Ask When Looking at a Model
The five-dimension model is the lens. This section is the lens applied as a checklist — the actual sequence of questions worth running, in order, the first time you're handed a model and asked to make it fast, small, or both. Organize by category, because each category has a natural owner (a model card, a profiler, a datasheet, a runtime's execution provider list) and asking the wrong category's question of the wrong source wastes time.
Model
- How many parameters, and how many FLOPs does one forward pass actually require at the batch size you'll actually run — not the batch size the paper reported?
- What are the tensor shapes at every stage, and do any of them collapse to degenerate cases (batch size 1, a spatial dimension of 1) where a kernel written for the general case will badly underperform a specialized one?
- Which operators dominate parameter count versus which dominate FLOPs — these are frequently not the same operators, and conflating them leads to optimizing the wrong one.
Memory
- What is the peak memory footprint across the whole forward pass, and does it fit in the target's SRAM, or does it spill to DRAM — and if it spills, at which layer does the spill first happen?
- How large are the weights alone, separate from activations, and does that number change under the precision you're planning to deploy at?
- What is the largest single intermediate tensor, and does that tensor alone already exceed a memory-constrained target's budget regardless of how well everything else is planned?
Compute
- Which operator type dominates total FLOPs — GEMM, convolution, attention, or elementwise — because the answer determines which entire lesson's toolkit is even relevant?
- Is the dominant operator compute-bound or memory-bound at its actual shape, per the Roofline test from Lesson 1, not per intuition?
- Are there elementwise or normalization operators that look cheap in FLOPs but are actually expensive in wall-clock time because they're memory-bound and un-fused?
Precision
- What precision does each operator actually need to hold accuracy, rather than what precision the framework happens to export by default — FP32, FP16, BF16, int8, int4 can coexist in one model?
- Is the precision choice per-tensor or per-channel, and does that choice match where the dynamic range actually varies (per-channel almost always wins for weights, per-tensor is usually adequate for activations)?
- Has the calibration or quantization-aware-training range actually been validated against a representative distribution, or is it a default that happens not to have crashed yet?
Hardware
- What compute units does the target actually expose — scalar core, SIMD/vector unit, GPU, NPU, DSP — and which of the model's dominant operators map cleanly onto which unit?
- Does the target have a tensor-core-style or systolic-array-style unit that only activates above a minimum matrix dimension, meaning small shapes silently fall back to a slower path?
- What is the actual memory bandwidth between the compute unit and the nearest fast memory tier, and does the model's arithmetic intensity clear that bar?
Runtime
- Which operators in this exact graph are actually supported by the target execution provider, and which ones silently fall back to a slower CPU implementation — a fallback that is often invisible until you profile for it?
- Are there unnecessary copies at the boundary between execution providers, or between the runtime and the application, that a zero-copy or memory-mapped path could eliminate?
- Does the runtime's memory planner get to see the real input shapes ahead of time, or is it forced into a slower, more conservative dynamic-shape path?
Compiler
- Can adjacent operators fuse, and if the compiler isn't fusing them, is that a missing pattern in the rewrite rules or a genuine correctness constraint (aliasing, side effects) blocking it?
- Can the graph's layout be normalized so only the minimum number of transforms — ideally one in, one out — are needed for the whole network, instead of one per mismatched operator boundary?
- Can constants fold and dead nodes eliminate before the graph ever reaches the runtime, so none of that work costs anything at inference time?
Performance
- What is the actual measured latency distribution — median, P95, P99 — not just the average, since a scheduler-, thermal-, or memory-pressure-induced tail can dominate user experience even when the mean looks fine?
- What is the achievable throughput at the batch size the serving system will actually use, and where does that number sit relative to the Roofline ceiling for the dominant operator?
- What is the energy per inference, and does that number — not latency alone — end up being the actual binding constraint on a battery-powered target?
6. The Complete Mental Picture
Everything above should compress into a single reflex: look at a model, and mentally expand it, top to bottom, through every layer this series built.
Run a transformer through this and it reads: QKV projections are GEMMs, GEMMs need a layout and a tile size, attention's softmax and KV cache change the shape of the memory-bandwidth problem between prefill and decode, every one of those pieces gets a precision assignment, kernel selection picks which concrete loop nest implements each GEMM, and the runtime decides which processor executes which piece and in what order.
Run a CNN through the same picture and it reads: convolution becomes FLOPs and MAC operations, tensor layout decides whether SRAM reuse is cheap or expensive, tiling decides how much of that reuse actually happens, int8 and SIMD or an NPU's MAC array execute the tiled kernel, and fusion decides how much of the resulting memory traffic between layers never has to touch DRAM at all.
That transformation — a model's name turning, unprompted, into this entire stack — is the actual finish line of this series.
Everything before this post was building the vocabulary for each box. This post is the claim that the boxes were never separate topics; they were always one picture.
7. The Recommended Learning Path — Seven Levels, and Why Each Is a Prerequisite
A learning path is only useful if it's also an argument about why the ordering is forced, not just a reading list. Each level below is a genuine prerequisite for the one after it — attempting the next level without it produces exactly the kind of cargo-cult understanding this whole series has tried to avoid, where a technique is applied because it's "supposed to help" rather than because its underlying constraint has actually been derived.
Level 1 — Mathematical foundations
│ (you cannot reason about arithmetic intensity
│ without first knowing what the arithmetic is)
▼
Level 2 — Computer architecture
│ (you cannot reason about why a memory-bound
│ kernel is slow without a memory hierarchy)
▼
Level 3 — Kernel engineering
│ (you cannot reason about a compiler's kernel-
│ selection decision without having hand-written
│ the kernels it's choosing between)
▼
Level 4 — ML-specific optimization
│ (you cannot reason about what a graph compiler
│ pass is *for* without knowing what quantization,
│ pruning, and fusion are trying to buy)
▼
Level 5 — Runtime
│ (you cannot reason about a compiler's lowering
│ target without knowing what the runtime actually
│ needs from the code it hands off)
▼
Level 6 — Compiler
│ (you cannot reason about hardware dataflow
│ architecture without first knowing what a
│ compiler is trying to map onto it)
▼
Level 7 — HardwareLevel 1, mathematical foundations — matrix multiplication, convolution, tensor shapes, attention, FLOP counting, arithmetic intensity, the quantization mapping — is the floor, because every later level's vocabulary assumes you can already read 2 × M × N × K off a GEMM shape without deriving it from scratch each time.
Lesson 1's function view and Lesson 3's affine quantization math are this level, in this series.
Level 2, computer architecture — cache hierarchy, SRAM versus DRAM, cache lines, memory bandwidth, SIMD, branch prediction, DMA, the Roofline model — is next because Level 1's arithmetic is meaningless as a performance argument until you know what it costs to feed that arithmetic.
A FLOP count with no memory-hierarchy context tells you nothing about wall-clock time; that's Lesson 1, Part 2's entire argument, and it's also most of the sibling Computer Architecture series this site runs alongside this one, which this series leans on directly for cache-line addressing and stride derivations.
Level 3, kernel engineering — actually implementing and optimizing GEMM, convolution, reduction, softmax, layer normalization with blocking, SIMD, cache locality, threading, and prefetching — is where Levels 1 and 2 stop being facts you can recite and become facts your hands have proven.
This is non-negotiable before Level 4, because every ML-specific optimization in Level 4 is ultimately justified by its effect on a kernel exactly like the ones this level makes you hand-write. Without having written one, "operator fusion avoids a memory round trip" is a sentence you're trusting rather than one you've watched be true.
Lesson 4's nine dimensions of kernel optimization and its worked GEMM and im2col derivations are this level.
Level 4, ML-specific optimization — quantization, pruning, sparsity, distillation, low-rank factorization, operator fusion, layout transformation — is where the kernel-level fluency from Level 3 gets pointed at model-level and graph-level decisions instead of a single loop nest.
You cannot evaluate whether fusing two operators is worth it without Level 3's understanding of what a memory round trip actually costs in cycles. You cannot evaluate whether int8 quantization will actually speed anything up without Level 3's understanding of which kernels are compute-bound versus memory-bound at their real shapes.
Lessons 2, 3, and 8 are this level.
Level 5, runtime — memory planning, tensor lifetime, buffer reuse, execution graphs, scheduling, asynchronous execution, device placement, zero-copy, DMA — comes next because it's the layer that takes everything Level 4 decided to do to the graph and actually executes it correctly and efficiently at inference time.
You cannot design a sound memory plan without Level 4's fused, laid-out, quantized graph already in hand — memory planning operates on the output of graph-level optimization, not before it.
Lesson 5, and the serving-specific runtime concerns of Lesson 6, are this level.
Level 6, compiler — ONNX, MLIR, TOSA, Linalg, LLVM, dialect conversion, pattern rewriting, bufferization, tiling, fusion, lowering — sits above the runtime in this ordering because a real compiler's job is to generate everything Levels 4 and 5 described by hand, automatically, across arbitrary graphs.
You cannot appreciate why progressive lowering through a dialect stack is the right architecture without having already felt, from Levels 3 through 5, the actual tension it resolves: needing to reason about both framework-level graph structure and hardware-level loop and memory detail, at different points, without either one contaminating the other.
Lesson 7 is this level, and MLIR's own documentation is the strongest primary reference for continuing past what this series covered.
Level 7, hardware — ARM Cortex and NEON, ARM's dot-product instructions, GPU architecture and Tensor Cores, NPU architectures, systolic arrays, MAC arrays, SRAM banking, DMA, accelerator dataflows — is last, not because it's least important, but because every one of the six levels before it is what makes a piece of specialized silicon legible instead of a black box.
A systolic array is, at the physical level, exactly Level 3's register-blocked GEMM micro-kernel and Level 2's cache-hierarchy reasoning, wired directly into silicon instead of expressed in software. A compiler targeting it, per Level 6, has to understand its dataflow well enough to generate code for it the way Level 6 generates code for a general-purpose core.
Lesson 10's latency budgets, energy accounting, and embedded-inference constraints are this level made concrete for the smallest end of the hardware spectrum; NPU dataflow architecture is this level made concrete for the specialized end.
Two lessons in this series don't sit cleanly inside a single level, and it's worth saying so rather than forcing a fit.
Lesson 9's profiling and benchmarking discipline is not a level at all — it's the feedback loop that validates every level's decisions against reality, and it belongs everywhere, continuously, not once in a sequence.
Lesson 6's serving-workload material — batching, latency-versus-throughput, the KV cache — is Level 5's runtime concerns applied to one specific, economically important production shape, and is worth revisiting once you've built real runtime fluency rather than treating it as a level of its own.
8. Six Projects That Will Actually Teach You This
Reading builds vocabulary. These six projects build the kind of understanding that survives being asked a hard question in an interview or a design review — because you will have already hit the wall the question is testing for.
| Project | What you build | Core skill exercised |
|---|---|---|
| 1. INT8 GEMM | FP32 GEMM, then int8 GEMM, then a SIMD int8 GEMM, benchmarked against each other | Levels 1–3: arithmetic, cache behavior, quantization mechanics, SIMD |
| 2. Tiny inference runtime | Tensor, Operator, Graph, Memory Arena, and Scheduler classes supporting MatMul, Add, ReLU, Conv, plus graph fusion | Levels 4–5: graph representation, memory reuse, scheduling |
| 3. Static memory planner | A tool that takes tensor lifetimes and assigns buffer offsets minimizing peak memory | Level 5, deep dive: the interference-graph-coloring problem worked end to end |
| 4. ONNX optimizer | A pass pipeline over a real ONNX graph implementing constant folding, dead-node elimination, and Conv+BatchNorm / Conv+ReLU fusion | Level 4 and Level 6: graph-level compiler passes on a production IR |
| 5. MLIR compiler | A toy ML graph lowered through MLIR into Linalg, into tiled loops, into LLVM, with the generated code actually inspected | Level 6, deep dive: progressive lowering, seen with your own eyes instead of read about |
| 6. Embedded inference | A small CNN taken from FP32 through int8, run on ARM with NEON and DMA-managed SRAM tiling, on real or emulated hardware | Level 7: the full stack compressed onto genuinely constrained silicon |
Project 1 — INT8 GEMM is the one to start with, and it is worth walking through in enough detail to actually attempt this week rather than someday.
Start small and square: 64×64×64, 128×128×128, 256×256×256, doubling up to 1024×1024×1024, so you can watch behavior change as the working set stops fitting in successive cache tiers.
Then add one shape that isn't square and isn't a toy — something in the neighborhood of a transformer feed-forward layer, M=128, K=768, N=3072 — because square benchmarks flatter naive code in ways that don't survive contact with a real model's actual shapes.
Write the naive triple-nested-loop FP32 version first and treat it as a correctness oracle, not a performance baseline. Then write an ikj-ordered version and confirm the cache-line-count argument from Lesson 4 by actually measuring, not assuming.
Convert to int8 next, accumulating into int32, and only after that add SIMD — NEON's dot-product instructions if you're on ARM, AVX2's _mm256_maddubs_epi16-family intrinsics if you're on x86.
Benchmark all of it against a real reference — OpenBLAS or Eigen for the FP32 ceiling, so you have an honest sense of how far "correct" is from "actually fast." Measure GFLOP/s, wall-clock latency, and, if your platform exposes it, cache-miss counts through perf or an equivalent.
The "aha" moment on this project tends to arrive twice, in opposite directions, and both arrivals are the actual point.
The first is disappointment: run your int8 SIMD kernel on the small, untiled shapes and it will often barely beat the naive FP32 loop, sometimes not beat it at all. At small K, the kernel is memory-bound, not compute-bound, and quadrupling the arithmetic throughput does nothing for a kernel that was never arithmetic-bound in the first place. That is Lesson 1's Roofline argument, no longer something you read, but something you just personally reproduced by accident.
The second arrival is the payoff: add the blocking Lesson 4 derived — tiling M, N, and K so a sub-block survives in L1 or L2 across all of its reuses — and watch the same int8 SIMD kernel jump to a genuine three-to-four-times speedup over the FP32 baseline on the larger shapes.
The gap between those two runs, on the same kernel, is the entire content of this series compressed into one measurement you took yourself.
The other five projects deserve the same seriousness even in brief.
The tiny inference runtime forces you to confront, in code, the exact distinction Lesson 5 spent an entire lesson establishing — that "what does this graph compute" and "where does every intermediate tensor live and when does each operator run" are different questions requiring different data structures, and that graph fusion is a real pass you have to write, not a property that falls out of having operators.
The static memory planner, built in isolation from a runtime around it, is the cleanest way to internalize that memory planning is genuinely graph coloring. Implement the greedy-by-birth-order algorithm on interval graphs, confirm it's provably optimal there, then implement greedy-by-size for the non-uniform-buffer-size case and watch it stop being provably optimal — which is exactly the NP-hard dynamic-storage-allocation reality Lesson 5 named but didn't have to prove to you.
The ONNX optimizer puts you inside a real production IR instead of a toy one. Conv+BatchNorm fusion specifically is worth doing by hand at least once — deriving the folded convolution weights and bias algebraically before writing the pass — because it's one of the few transformations in this series simple enough to verify by hand and consequential enough to matter in every deployed CNN.
The MLIR compiler project is the single highest-leverage project for this reader specifically. Build the toy-graph-to-Linalg-to-tiled-loops-to-LLVM pipeline from Lesson 7, then actually read the generated IR at each stage — you will be looking at your own professional domain from the inference-application side instead of the compiler-internals side for the first time.
Embedded inference closes the loop by making every constraint from Lesson 10 physical — kilobytes instead of gigabytes, no operating system to fall back on, an allocator failure that's a hard fault rather than a slow path. It's the project most likely to surface a bug that only exists because you stopped assuming infinite memory.
Further Reading
- ONNX Runtime, "Performance" — the official documentation home for graph optimization, quantization, execution providers, and the arena/memory-pattern behavior referenced throughout Lesson 5 and this post.
- MLIR, "Code Documentation" — the primary reference for the dialect stack, pass infrastructure, and progressive-lowering philosophy this post's compiler-engineering section leans on directly.
- MLIR, "Quantization" — the dialect-level documentation for lowering floating-point computation into the integer representations Lesson 3 derived from first principles.
- NVIDIA, "Best Practices for TensorRT Performance" — a production reference for GPU inference optimization: batching, layer fusion, Tensor Core usage, and profiling methodology, from the vendor side of the stack this series studied conceptually.
- Hennessy, J. L. and Patterson, D. A., Computer Architecture: A Quantitative Approach, 6th ed., Morgan Kaufmann / Elsevier — the standard graduate reference for the memory-hierarchy and quantitative-performance reasoning underlying Lesson 1 and Lesson 2. Elsevier product page
- Patterson, D. A. and Hennessy, J. L., Computer Organization and Design, 6th ed., Morgan Kaufmann / Elsevier — the gentler companion volume, closer to this series' own fetch-decode-execute-to-kernel path. Elsevier product page
- Bryant, R. E. and O'Hallaron, D. R., Computer Systems: A Programmer's Perspective, 3rd ed. — the strongest single book for connecting C/C++ source directly to assembly, memory, and measured performance, which is exactly Section 2's argument made at book length. csapp.cs.cmu.edu
- Kirk, D. B. and Hwu, W. W., Programming Massively Parallel Processors: A Hands-on Approach, 4th ed., Morgan Kaufmann / Elsevier — the standard text for GPU execution models and parallel programming patterns behind Level 7's GPU material. Elsevier product page
- Sze, V., Chen, Y-H., Yang, T-J., and Emer, J. S., "Efficient Processing of Deep Neural Networks: A Tutorial and Survey" — the freely available survey covering hardware and software techniques for efficient DNN execution, spanning most of what Levels 4 through 7 touch on in one document. arXiv:1703.09039
Closing This Series
Nineteen parts ago, this series opened with a single claim: that running a trained model is an engineering discipline in its own right, not an afterthought tacked onto training.
Everything since has been that claim, defended one layer at a time — the Roofline model defended it at the level of arithmetic, quantization defended it at the level of bits, kernel engineering defended it at the level of loops and registers, memory planning and the runtime defended it at the level of buffers and schedules, the compiler defended it at the level of IR, and Lesson 10 defended it at the level of joules and milliseconds on real, unforgiving silicon.
None of that was ever a list of independent tricks. It was one argument, told from five different angles — compute, data movement, parallelism, precision, and scheduling — because those five angles are, as best this series can tell, the whole shape of the problem.
If you came to this series with a C++ systems background and an MLIR-based NPU compiler background, as Sections 2 and 3 assumed, then a good part of what you just read wasn't new information so much as new labels on instincts you already had.
That's not a coincidence, and it's not a coincidence worth being modest about: inference engineering, at its foundation, is systems programming and compiler engineering, applied to a workload whose mathematics happens to be linear algebra instead of general-purpose control flow.
The interference graph you build to plan tensor memory is the interference graph you'd build to allocate registers. The dialect stack you lower a graph through is the AST-to-machine-code pipeline you already know, stretched across a wider abstraction gap. The loop nest you'd hand-tune for a hot C++ path is the loop nest a kernel-selection pass is choosing on your behalf.
You were not learning a new field so much as watching a field you already understood get renamed, shape by shape, until the renaming stopped and you could see it was the same field the whole time.
What changes now is only what's asked of you.
Reading a Roofline diagram is not the same as pulling one out of a real profile on real hardware under real deadline pressure. Understanding graph coloring is not the same as watching your own memory planner fail to hit the theoretical minimum on an irregular graph and having to figure out why. Recognizing MLIR's dialect stack as familiar is not the same as writing the lowering pass that makes a new operator work on it.
That gap — between recognizing an idea and having built it — is exactly what Section 8's six projects exist to close, and it's the only gap left between finishing this series and actually owning the discipline it describes.
The five questions this whole series reduces to are worth carrying forward past the last line of this post, because they don't expire when the reading does: how much computation, how much data movement, how much parallelism, what precision, and when and where should it run.
Ask them of the next model you're handed, the next kernel you profile, the next compiler pass you write, the next chip you get a datasheet for.
That habit — not any single fact in these nineteen parts — is the actual thing this series was trying to hand you.