The hardware underneath: pipes, DRAM and clocks

1. What chapters 2–5 rounded off

Every model of a machine leaves things out; the useful question is whether what it leaves out ever changes a decision. Part I’s model left out three things that occasionally do:

Each section below follows the same pattern as chapter 4: a number from NVIDIA’s documentation, what it predicts, and an experiment in lab 4b that checks it on your card. Two animated figures show the two places where timing is the whole story: one quadrant’s pipes, cycle by cycle, and one memory channel, nanosecond by nanosecond.

2. Issue and execute: the pipes of a quadrant

Chapter 2’s scheduler picks one ready warp per cycle and issues its next instruction. Here is the part it skipped: where the instruction goes. NVIDIA’s diagram of the RTX Blackwell SM shows each of the four quadrants (Nsight calls them SMSPs) with its own L0 instruction cache, warp scheduler, dispatch unit and 64 KB register file, feeding these units:

• 32 FP32/INT32 units. On Blackwell each one can do either a float or an integer operation in a given cycle. Nsight Compute splits them into two pipes, fmaheavy and fmalite: both do FP32 add, multiply and FMA, but only the heavy half does integer multiply (IMAD). That is why IMAD runs at half the FP32 rate.
• One SFU (special function unit; Nsight’s xu pipe) for MUFU, the multi-function instruction: reciprocal, reciprocal square root, base-2 log and exponential, sine, cosine. 16 per SM per cycle, so 4 lanes per quadrant.
• FP64: 2 per SM per cycle on GeForce Blackwell. Half a lane per quadrant.
• 4 load/store units that send memory instructions, shuffles and barrier operations to the L1 (section 6).
• A tensor core (chapter 8), and a uniform datapath for values every thread shares (section 6).

Two numbers describe a pipe. Its throughput is how often it can accept a new warp instruction: a 32-lane pipe takes one every cycle, a 16-lane pipe one every 2 cycles, the 4-lane SFU one every 8. Its latency is how long until the result can be used, about 4 cycles for an FMA. A pipe is like the conveyor of chapter 3’s Little’s law: it can hold several instructions at once, at different stages.

NVIDIA publishes the throughputs as results per clock per SM in the CUDA C++ Best Practices Guide (the column for compute capability 12.0, which is your card). Dividing by 4 quadrants and then into 32 threads gives the cycles one warp instruction occupies its pipe:

The scheduler is a pipe too: one warp instruction per cycle per quadrant, whatever kind it is. That single fact has a consequence chapter 4 hinted at. FP32 peak needs an FFMA issued every cycle, so every other instruction in the loop, an address computation, a load, a loop counter, a comparison, takes an issue slot that an FFMA could have used. Peak FLOP/s is a statement about a loop made only of FMAs. Real loops get close only when the FMAs vastly outnumber everything else, which is what tiling buys in chapter 7: each loaded value is reused in many FMAs.

3. Latency: waiting on the compiler, waiting on a scoreboard

An instruction that needs the result of an earlier one must wait for it. How the hardware knows how long to wait depends on the pipe.

Fixed latency. FP32, integer and logic instructions always take the same number of cycles; the Nsight Compute Profiling Guide states that “FFMA-to-FFMA dependency latency is 4 cycles”. Because it is fixed, the hardware does not track it at all. The compiler does: every SASS instruction carries a few control bits, set by the compiler, that tell the scheduler how many cycles to hold this warp before its next instruction. (NVIDIA doesn’t document their format; Jia et al. reverse-engineered it, see section 9.) A warp held this way shows up in Nsight Compute as the stall reason wait: “waiting on a fixed latency execution dependency”.

Variable latency. A shared-memory load might hit a bank conflict; a global load might hit L1, L2 or DRAM. The compiler cannot know, so the hardware keeps a small set of scoreboards per warp: the instruction that produces a value sets a scoreboard, the one that consumes it waits until the scoreboard clears. There are two kinds, and they are the two scoreboard stalls you met in chapter 5:

• short_scoreboard: waiting on an MIO (memory input/output) operation that does not go to L1TEX: shared memory above all, but also MUFU and dynamic branches. Latency in the tens of cycles.
• long_scoreboard: waiting on L1TEX: global, local and texture memory. Hundreds of cycles if it goes to DRAM.

Latency is hidden the way chapter 3 hid DRAM latency, by Little’s law: work in flight = latency × throughput. For the FMA pipe that is 4 cycles × 1 warp instruction per cycle = 4 independent FMAs in flight per quadrant. They can come from 4 warps that each have one chain (thread-level parallelism, TLP) or from one warp with 4 independent chains (instruction-level parallelism, ILP), or any mix. Chapter 4’s compute benchmark used four chains per thread for exactly this reason.

Reading Fig. 1 against chapter 5 gives a map from stall reason to hardware:

4. The SFU, and the price of exp

When you write expf(x) in a kernel, the compiler does not emit one instruction. The SFU only computes 2^x (MUFU.EX2), approximately. To build an accurate expf, the compiler multiplies by log₂ e, splits the result into an integer and a fraction, fixes up the rounding with a few FMAs, calls the SFU, and handles overflow and underflow. CUDA offers two versions, and the Programming Guide lists their accuracy:

The same split exists for __sinf, __logf, __fdividef and others. The compiler flag -use_fast_math replaces every call with its intrinsic in one go. It implies --ftz=true (denormals flushed to zero), --prec-div=false and --prec-sqrt=false (fast, approximate division and square root) and --fmad=true. That is a lot of semantics to change with one flag; in this course we write intrinsics by hand where we want them and leave the rest accurate.

How much does it matter? Two cases, one where it doesn’t and one where it does.

Softmax on your card: not at all. A softmax over FP32 values reads 4 bytes and writes 4 bytes per element and does one exp. At the practical 590 GB/s that is about 74 billion elements per second. The SFUs can do 16 × 48 SMs × 2.51 GHz ≈ 1.9 trillion MUFU per second, 26× more. Softmax is memory-bound by a wide margin; whether its exp takes one MUFU or ten instructions doesn’t show up in its time. Chapter 4’s rule holds: on the roofline, removing work you don’t wait for buys nothing.

Attention: it depends on the card. In FlashAttention (Part III) the scores never touch DRAM. Each score costs about 4d FLOPs on the tensor cores (a dot product of length d to make it, and its share of the product with V) and one exp on the SFU, and both units work on the same tile at the same time. Compare their speeds:

Tensor cores got 16× faster per SM between the two chips; the SFU did not. On a B200, attention with head dimension 128 spends as long on exp as on matrix multiplication, and the two can overlap at best. FlashAttention-4 (2026) responds by computing part of the exponentials without the SFU: a cubic polynomial evaluated with FMAs on the FP32 pipe, which has 8× the SFU’s throughput and is otherwise idle. On your card that trick isn’t needed. The lesson generalizes: a unit that didn’t scale becomes the bottleneck when everything around it did, and the fix is often to move its work to a unit that did.

The fourth and fifth acts of Fig. 1 show the same thing at the scale of one quadrant: one MUFU per instruction group saturates the SFU at one warp instruction per 8 cycles, but up to seven other instructions can issue in its shadow for free.

5. FP64 and the accidental double

GeForce cards are built for graphics and AI, which don’t need 64-bit floats, so NVIDIA gives them just enough FP64 to run code that uses it: 2 operations per clock per SM, 1/64 of FP32. (Datacenter chips like the B200 have far more.) That is fine until a kernel uses double without meaning to. The classic way:

__global__ void scale(float* y, const float* x, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) y[i] = x[i] * 0.5;    // 0.5 is a double literal
}

C++ promotes x[i] to double, multiplies, and converts back. In SASS that is F2F.F64.F32, DMUL, F2F.F32.F64: three instructions at the FP64 rate, 64 cycles each per warp. The fix is one character, 0.5f. Other ways in: pow(x, 2.0), sqrt(2.0) used as a constant, a double constant in a header, an unsuffixed literal like 1e-6 in a comparison. The sixth act of Fig. 1 shows what it looks like: the issue slot idles 98% of the time while the FP64 unit chews through one warp instruction every 64 cycles.

Will this slow a memory-bound kernel like scale? Each element needs three FP64 operations, and the GPU does 2 × 48 × 2.51 GHz ≈ 240 billion per second: 80 billion elements per second, about the same as the 74 billion that DRAM allows. So the accident turns a cleanly memory-bound kernel into one limited by both, and any kernel with more than one stray double per element becomes FP64-bound. Two ways to catch it, both cheap:

cuobjdump -sass ./prog | grep -E "DMUL|DADD|DFMA|F2F\.F64"      # any FP64 instruction in the binary
nvcc -Xptxas --warn-on-double-precision-use ...                   # ptxas warns at compile time

The second is worth adding to every Makefile in this course that has no business using double. In Nsight Compute, the Compute Workload Analysis section shows the fp64 pipe’s utilization; anything above zero in a float kernel is an accident.

6. The rest of the quadrant

A few more units appear in Nsight Compute’s pipe list. You don’t need to tune for them, but you will see their names:

• Uniform datapath (uniform): a scalar unit for values that are the same in every thread of the warp: base pointers, loop counters, blockIdx-derived offsets. It has its own small register file (uniform registers, UR in SASS). Work on it doesn’t take an FP32/INT32 lane, though each instruction still uses the issue slot.
• Load/store (lsu): 4 units per quadrant. Every global, local and shared memory instruction passes through here on its way to the L1, and so do shuffles, special-register reads (S2R, e.g. reading threadIdx) and block-level barrier operations. That is why shuffles show up as memory-like traffic in the profiler.
• Branch units: the convergence barrier unit (cbu) handles warp-level convergence and branches, the divergence of chapter 2; the address divergence unit (adu) handles indirect jumps and also serves constant loads.
• Texture units (tex): 4 per SM, for graphics texture sampling. CUDA uses them for some read-only loads.
• Tensor core (tc): chapter 8.

In the Compute Workload Analysis section, Nsight Compute shows the utilization of every one of these pipes. When the FMA pipe is not the busiest, that chart tells you which one is.

7. Below the L2

Chapter 3 stopped at the L2: on a miss, a 32-byte sector comes from “DRAM”. Here is the path it takes, from top to bottom:

1. L2 slices. The 48 MB L2 is not one cache but many slices; each physical address belongs to exactly one slice, picked by a hash of its bits. NVIDIA doesn’t publish the slice count for your chip; Nsight Compute’s lts__ metrics (L2 “tag stage”) count per slice.
2. A crossbar connects every SM to every slice, so any SM can reach any address at the same cost.
3. Six memory controllers, each 32 bits wide, each attached to one 2 GB GDDR7 chip: 192 bits, 12 GB.
4. Four channels per chip. A GDDR7 chip is split into four independent 8-bit channels, so the card has 24. Each pin runs at 28 Gbit/s, so each channel moves 28 GB/s, and 24 × 28 = 672 GB/s, the number from chapter 1. GDDR7 sends 1.5 bits per symbol by using three voltage levels (PAM3) instead of two; that is how it reaches 28 Gbit/s per pin.
5. Bursts. One access moves 32 bits per pin over 8 pins: 32 bytes, exactly one L2 sector, in about 1.14 ns.

The address hash matters more than it looks. Early GPUs picked the channel from a few plain address bits, so a strided access pattern, like walking down a matrix column, could land on one channel while the others sat idle (“partition camping”). Today’s hash scatters any regular pattern over all channels and slices; you no longer need to think about it.

Banks and rows

Inside each channel the memory is divided into banks, and each bank into rows of a few KB. A bank can only read from one row at a time, through its row buffer:

• ACTIVATE copies a whole row into the row buffer. That takes about 10–15 ns.
• READ or WRITE moves one 32-byte burst between the row buffer and the pins. As long as the row stays open, these can go back to back at the full 1.14 ns per burst.
• PRECHARGE writes the row back and closes it, another 10–15 ns, before the bank can open a different row.

On top of that, DRAM limits how many rows may be activated in any short window, to bound the current the chip draws (DDR’s rule is four; the model below uses the same). A stream of reads that keeps hitting open rows runs at full speed. A stream where every burst needs its own row pays an ACTIVATE and a PRECHARGE per 32 bytes and hits the activation limit. Fig. 2 runs one channel through five traffic patterns.

The model’s numbers over the 110 ns window:

This is where the practical ceiling of chapter 4 comes from. Even a perfect stream loses a few percent to refresh (every row must be rewritten every few tens of milliseconds, and the chip pauses for it) and to read/write turnaround (a kernel that both reads and writes makes the bus switch direction, and each switch costs idle nanoseconds). Hence chapter 4’s 590 GB/s, 88% of 672.

And this is the hardware reason behind a rule you will meet again in Part III: random access is fine when the unit of randomness is large. Chapter 3 already showed that a warp should use all 32 bytes of each sector it touches. Below the L2 the rule extends: scattered 32-byte sectors waste most of DRAM’s time on opening rows, while scattered chunks of a few KB stream almost as well as one contiguous array. An embedding table of 4 KB rows, or a KV cache stored in pages of several KB (Part III’s paged attention), can be read in random order at nearly full bandwidth.

One more rung: the TLB

Every address a kernel uses is virtual. Translating it to a physical address goes through a page table, cached in small TLBs (translation lookaside buffers) at each level. GPU memory is mapped in large pages, typically 2 MB, so a TLB with a few thousand entries covers gigabytes. But not unlimited gigabytes: chapter 3’s latency staircase can creep up past a few hundred MB, and accesses spread randomly over the whole 12 GB miss the TLB and pay a page-table walk on top of DRAM latency. NVIDIA doesn’t publish the TLB sizes; Jia et al. measured them on earlier GPUs by pointer-chasing with one access per page, and experiment H in lab 4b does the same on yours.

8. Clocks and power

Chapter 4 computed the peak as 6,144 lanes × 2 FLOPs × 2.51 GHz = 30.9 TFLOP/s. The 2.51 GHz is the rated boost clock: a promise that the card will reach it under typical load. The clock it actually runs at is chosen continuously by the card’s firmware (GPU Boost) as the highest frequency allowed by three limits: the power budget (250 W total graphics power on the 5070), the temperature, and the voltage. A cool card running light work usually goes above the rated boost; a card running heavy tensor-core work at its power limit may go below. RTX Blackwell changes its clock far faster than earlier GPUs (NVIDIA says about 1000×), so the clock can now move within a single frame, and so within a long kernel.

The memory clock is separate and much steadier: GDDR7 runs at its rated 28 Gbit/s whatever the SM clock does. That splits kernels into two kinds:

Three practical consequences:

• The ridge point moves with the clock. At 590 GB/s it is 30.9 TFLOP/s / 590 GB/s ≈ 52 FLOPs per byte at 2.51 GHz, about 42 at 2.0 GHz and 59 at 2.85 GHz. A kernel near the ridge can switch between compute- and memory-bound as the card heats up.
• Report the clock with every compute-bound measurement. A 10% speedup that coincides with a 10% higher clock is not a speedup. Warm the card up before timing, and log the clock while you measure:

  nvidia-smi --query-gpu=clocks.sm,clocks.mem,power.draw,temperature.gpu --format=csv -lms 200
  nvidia-smi -q -d PERFORMANCE      # why the clock is what it is: power cap, thermal, idle...

• Nsight Compute fixes the clock (chapter 5) so that its results are comparable between runs. That is why its durations can differ from your benchmark: they are measured at a different clock.

Lab 4b measures the clock directly: a kernel reads the SM’s cycle counter (clock64()) at its start and end, and dividing cycles by nanoseconds gives the real frequency during that kernel.

9. The rest of the chip, and where to read more

Outside the SMs and the memory system, a few more units take part in every CUDA program:

• The GigaThread engine is the block scheduler of chapter 2: it hands blocks to SMs as space frees up.
• Copy engines are DMA units that move data between CPU memory and GPU memory over PCIe without using any SM. They are what lets cudaMemcpyAsync overlap with running kernels (PMPP ch. 23).
• NVENC and NVDEC encode and decode video, and the display engine drives the monitor; none of them matters for CUDA, except that the desktop shares your GPU (chapter 5).
• The PCIe 5.0 ×16 interface connects all of it to the CPU (chapter 2’s film).

Where to go deeper, once Part II is done:

10. Lab 4b: measure the hardware

cd ~/unibo-lessons/mlsys/labs/lab04b-pipes-and-dram
make test          # quick correctness checks (every kernel computes what it claims)
make sass          # count the instructions in each loop: is it the one it claims to be?
make bench-pipes   # A: latencies, B: throughput per instruction, C: warps x ILP, D: MUFU + k FFMA
make bench-dram    # E: read/write/copy, F: working-set sweep, G: random chunks, H: TLB reach

Check your understanding

Sources and further reading

• NVIDIA, NVIDIA RTX Blackwell GPU Architecture (whitepaper, 2025): the SM diagram (per quadrant: L0 instruction cache, warp scheduler, dispatch at 32 threads per clock, 16,384 × 32-bit registers, 32 FP32/INT32 cores, 4 load/store units, an SFU, a tensor core; per SM: 128 KB L1/shared, 4 texture units, 2 FP64 cores), the GB205 specifications (48 SMs, 2,512 MHz boost, 48 MB L2, 192-bit GDDR7 at 28 Gbps, 672 GB/s, 250 W), and accelerated frequency switching.
• NVIDIA, CUDA C++ Best Practices Guide, “Arithmetic Instructions”: results per clock cycle per multiprocessor by compute capability (the table in section 2 uses the 12.0 column).
• NVIDIA, CUDA C++ Programming Guide, “Mathematical Functions” and “Intrinsic Functions”: maximum ulp errors of expf and __expf; the effect of -use_fast_math. NVCC documentation for --warn-on-double-precision-use.
• NVIDIA, Nsight Compute Profiling Guide: pipeline descriptions (fmaheavy, fmalite, xu, lsu, uniform, cbu, adu, fp64), the 4-cycle FFMA dependency latency, and the warp stall reasons quoted in section 3.
• JEDEC, JESD239 GDDR7 SGRAM standard, and vendor summaries of it: four independent channels per device, PAM3 signalling. Bank counts, row sizes and timings in Fig. 2 are illustrative, not taken from the standard.
• Zhe Jia, Marco Maggioni, Benjamin Staiger, Daniele P. Scarpazza, Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking (2018), and Jia, Maggioni, Smith, Scarpazza, Dissecting the NVIDIA Turing T4 GPU via Microbenchmarking (2019): latencies, throughputs, TLB measurements, and the methods lab 4b borrows.
• FlashAttention-4 (2026), arXiv:2603.05451: the exponential as a bottleneck on Blackwell datacenter GPUs and its polynomial emulation on the FMA units.
• Hwu, Kirk, El Hajj, Programming Massively Parallel Processors, 5th ed.: ch. 6 (DRAM), section 9.3 (atomics), appendix C (virtual memory), ch. 23 (streams).