3. Pipelining, multi-buffering, latency hiding — ML Hardware & Low-Latency Systems

The core mental model

Two numbers describe any pipeline and they are independent. Initiation interval ( $II$ ) is the gap between accepting successive inputs; throughput is $1/II$ . Depth is the time from an input entering to its own output emerging. A design with $II = 1$ and depth 200 is superb at throughput and useless against a 50-cycle deadline. Most engineering instinct optimises $II$ , because most work is throughput work, and the instinct is so ingrained that people report depth improvements as regressions when they cost a cycle of $II$ .

The same pipeline read two ways. Initiation interval is the horizontal gap between successive inputs — how often you can start one. Depth is the distance from an input entering to its own answer leaving. Deepening the pipeline shortens II and lengthens depth, which is a win for throughput work and a loss for a deadline, and the two are routinely reported as if they were one number.

The cost of running $N$ items through a pipeline is $\text{depth} + N \cdot II$ , so efficiency is $\frac{N \cdot II}{\text{depth} + N \cdot II}$ . That fraction is the whole subject. At large $N$ it tends to 1 and depth is amortised into irrelevance; at $N = 1$ it is $II/(\text{depth} + II)$ and you pay the entire depth for one answer. Every latency-hiding technique is a way of finding other work to occupy the depth: warp-level parallelism on a GPU (the scheduler swaps in another warp on a stall), the reorder window and prefetchers on a CPU, explicit software pipelining on an FPGA. They differ in who does the finding — hardware, compiler, or you — and not at all in what they do.

Multi-buffering is the explicit form. While computing on buffer $i$ you load buffer $i+1$ , which costs 2× the storage and hides one stage of memory latency behind one stage of compute. If the load is slower than the compute you need more buffers, and the count is a ratio, not a taste: prefetch distance $D = \lceil L_{\text{mem}} / T_{\text{compute}} \rceil$ . Hopper’s cp.async and TMA exist to make this cheap — moving global to shared without staging through registers, with an explicit commit-and-wait group model that is software pipelining exposed as an instruction set. And then the reframe that Part B is built on: at batch 1 against a deadline there is no other work to overlap with, every technique in this module returns exactly nothing, and the only remaining lever is reducing depth itself.

Numbers worth memorizing

Quantity	Value	Note
Pipeline cost for $N$ items	$\text{depth} + N \cdot II$	the formula the whole module hangs on
Efficiency at $N = 1$	$II / (\text{depth} + II)$	pipelining buys nothing
Prefetch distance	$\lceil L_{\text{mem}} / T_{\text{iter}} \rceil$	derive it, do not guess
HBM latency in GPU cycles	~400 ns ≈ 700 cyc at 1.75 GHz	what you are hiding
Warps needed for bytes in flight	~80 sectors/SM (Module 1)	occupancy alone rarely reaches it
Max warps / SM	64 (H100)	so you also need ILP, not just occupancy
Double-buffer shared memory cost	2 × tile bytes	direct hit to occupancy
CPU L2 prefetch distance, streaming	8–16 lines (512–1024 B)	prefetchers stop at 4 KB page bounds
Branch mispredict penalty	~15–20 cyc	≈ pipeline depth
FPGA target	$II = 1$ , depth as small as timing allows	the discipline inverts

The retiming trade, which is the thing most people have not internalised:

T_{\text{total}} = \text{depth} \times T_{\text{clk}}

Adding pipeline registers shortens the critical combinational path, so $T_{\text{clk}}$ falls and $F_{\max}$ rises — but depth grows. Minimise the product, not either factor. A 10-stage design at 500 MHz (20 ns) beats a 4-stage design at 150 MHz (26.7 ns), and a 40-stage design at 700 MHz (57 ns) loses to both.

Critical thinking

How many warps do you actually need to hide HBM latency? Work it out rather than quoting a rule of thumb.

Little’s Law, in bytes. To sustain 3.35 TB/s at ~400 ns you need about 1.34 MB in flight, which across 132 SMs is roughly 10 KB per SM, or about 80 outstanding 128 B sectors per SM.

Now count what the warps supply. If each warp has one coalesced 128 B load outstanding, 64 warps give you 64 sectors — short of 80, and that is at maximum occupancy, which you will not have because your tiles consume shared memory and registers. At a realistic 50% occupancy you are at 32 sectors, less than half of what the memory system needs.

The conclusion is the useful part: occupancy alone does not get you there, and instruction-level parallelism is not optional. Unroll so each thread has 2–4 independent loads in flight before it consumes any of them. This is why a kernel at 30% occupancy with deep unrolling routinely beats the same kernel at 75% occupancy without it, and why “raise occupancy” is bad advice stated as a goal rather than as a means to a bytes-in-flight target.

You add double buffering to a tiled kernel and get no speedup at all.

Four candidates, and they are distinguishable by measurement rather than by argument:

You were never memory-latency bound. Warp-level parallelism was already covering the load latency, so there was nothing left to hide. Check stall reasons: if the dominant stall was not long_scoreboard (waiting on memory), double buffering had no target.
The extra shared memory cost you a block per SM. You bought overlap and paid in occupancy, and the payment exceeded the purchase. Compute the occupancy before and after rather than assuming.
The compiler already did it. With enough registers and a simple loop, nvcc will software pipeline on its own. Your explicit version may be the same code with more source.
Your loads are not actually asynchronous. A __syncthreads() between the issue and the use collapses the overlap, and so does consuming the loaded value too soon. Double buffering is a scheduling claim; if the dependence structure does not permit the schedule, writing two buffers changes nothing.

A design has II = 1 and depth 200 cycles. Your deadline is 50 cycles. What do you actually do?

Attack depth, and accept whatever happens to $II$ as long as it stays at 1.

Rebalance, then retime. Depth × clock period is the objective. If the 200 stages exist to reach a high $F_{\max}$ you may be able to merge stages: fewer, longer stages at a lower clock can reduce total time even though the stage count improves.
Flatten the reductions. A sequential accumulation of $N$ terms is depth $N$ ; an adder tree is $\lceil \log_2 N \rceil$ . This is usually the single largest cut available.
Carry-save the arithmetic. Deferring carry propagation through a compressor tree and resolving it once at the end removes the carry chain from every intermediate add.
Speculate. Compute both sides of a decision in parallel and select late. You spend area to take a dependency off the critical path, which is Module 4’s entire thesis and Module 11’s central trick.
Move work off the path. Anything computable before the input arrives should be precomputed; anything needed only after the answer is emitted should be moved after it.

What you must not do is the reflex from throughput work — adding stages to raise $F_{\max}$ — the one thing that reliably makes this design worse.

What goes wrong if prefetch distance is too short, and what goes wrong if it is too long?

Too short and the data has not arrived when you need it, so you stall anyway and the prefetch bought nothing but instruction overhead. That failure is obvious.

Too long is the interesting one. The prefetched line sits in cache waiting to be used, and while it waits it can be evicted by the very stream that prefetched it — with $D$ too large and several streams running, your working set exceeds the cache and you thrash, which is strictly worse than not prefetching. You can also exhaust the outstanding-miss budget (10–16 line-fill buffers on a CPU core) so the prefetches themselves block demand misses.

So there are two bounds, not one. Set $D = \lceil L / T_{\text{iter}} \rceil$ for coverage, then check $D \times \text{line size} \times \text{streams} < \text{cache size}$ for survival. On a GPU the identical constraint is that the number of buffers times the tile size must fit in shared memory while leaving enough for the occupancy you need.

At batch 1 with a hard deadline, is there genuinely nothing to pipeline?

There is, and getting this right matters because the slogan overshoots.

You still pipeline for throughput across successive independent inputs. Market ticks keep arriving; $II = 1$ with depth 30 means you accept one input per cycle and each answer emerges 30 cycles later. Both properties are useful and they do not conflict.

What you cannot do is trade depth for $II$ , which is the trade every throughput-oriented technique offers. Deepening the pipeline to raise the clock, batching inputs to amortise a setup cost, adding a stage to simplify a critical path — all of these are normally free wins and all of them are now charged directly against the deadline.

So the discipline is: pipeline to exactly the depth timing closure requires, and not one stage more. That inverts the usual FPGA instinct, where you pipeline aggressively because $F_{\max}$ is the metric on the report. Here the report is wrong; total time is $\text{depth} \times T_{\text{clk}}$ and the deadline is on that product.

Self-check

Write the cost of running N items through a pipeline, and evaluate it at N = 1.

$\text{depth} + N \cdot II$ , with efficiency $\frac{N \cdot II}{\text{depth} + N \cdot II}$ .

At $N = 1$ that efficiency is $II/(\text{depth} + II)$ — for $II = 1$ and depth 30, about 3%. You pay the full depth for a single answer, which is the formal statement of why pipelining does nothing for a batch-of-one deadline.

Derive the prefetch distance for a loop whose iteration takes 5 ns against a 100 ns memory latency, then state the constraint that bounds it from above.

$D = \lceil 100/5 \rceil = 20$ iterations ahead.

The upper bound is capacity: $D \times \text{line size} \times \text{streams}$ must fit in the cache level you are prefetching into, or the prefetched lines are evicted before use. With 20 iterations, 64 B lines and 4 streams that is 5 KB, comfortable in L1; with 200 iterations and 16 streams it is 200 KB, which will thrash L1 and L2 and be slower than no prefetching at all.

Why is “more pipeline stages raises Fmax, so it is faster” wrong for a deadline?

Because the quantity that matters is $\text{depth} \times T_{\text{clk}}$ , and adding stages moves both factors in opposite directions. Going from 4 stages at 150 MHz (26.7 ns) to 10 at 500 MHz (20 ns) is a win; going on to 40 stages at 700 MHz (57 ns) is a large loss. $F_{\max}$ is a constraint you must satisfy, not the objective.

You have 64 warps per SM, each with one coalesced load outstanding. Is that enough to saturate H100 HBM?

No. You need about 1.34 MB in flight, roughly 80 sectors per SM; 64 warps × one 128 B load gives 64, and that assumes full occupancy, which real tiles do not achieve. The gap is closed with instruction-level parallelism — unrolling so each thread has several independent loads in flight — which is why deep unrolling at moderate occupancy often beats high occupancy without it.

Name three ways to cut pipeline depth that do not involve raising the clock.

Replace sequential accumulation with an adder tree ( $N \to \log_2 N$ ); use carry-save arithmetic so carries propagate once at the end rather than in every intermediate add; and speculate — compute both branches of a decision in parallel and select late, spending area to remove a dependency from the critical path. Merging over-pipelined stages and precomputing anything that does not depend on the input also belong on the list.

Double buffering produced no speedup. Give the first measurement you would take.

Stall reasons. If the dominant stall was not waiting on memory (long_scoreboard) there was nothing to hide and the technique had no target. The second measurement is achieved occupancy before and after, since the extra shared memory may have cost a block per SM — buying overlap and paying more for it than it was worth.

The core mental model#

Numbers worth memorizing#

Critical thinking#