Connect the modern system

Compute a tile with a matrix instruction

In Unit 5's tiled multiplication, each thread calculated products and added them to its output. A matrix instruction asks specialized hardware to perform these products and sums across a whole tile. It multiplies input matrices A and B and adds the result to an existing matrix C:

Here i selects an output row and j an output column. The sum over k pairs values across row i of A with values down column j of B. For example, D[0,1] in the explorer below is 0 + 1 × 6 + 2 × 8 = 22. Each A value contributes to several output columns, and each B value to several output rows. Matrix hardware can reuse those values while forming the tile.

Press Add next k once. At a fixed k, multiplying each value in column k of A by each value in row k of B produces a contribution to every output. The explorer adds all four contributions to D together.

Real matrix instructions accept particular tile shapes and input types. CUDA's warp matrix multiply-accumulate (WMMA) operations commonly support a 16 × 16 × 16 shape: M = 16 output rows, N = 16 output columns, and K = 16 products per output. Other supported shapes depend on the types and hardware. To use such a shape for a 2 × 2 problem, a program could pad the inputs with zeros. The extra products would occupy hardware without contributing useful work. The explorer's two steps describe the calculation; they do not specify hardware instructions or cycles.

The hardware also requires a particular layout, or arrangement of operands in memory and registers. A fragment is the portion of a matrix held by one participating thread; it need not be a complete row. The wmma family coordinates threads within a warp. The wgmma family, or warp-group matrix multiply-accumulate, coordinates four warps and has its own completion rules. A tuned matrix library provides a useful performance baseline because it handles these details. When comparing advertised peaks, check the input formats and the format of the running sums (accumulators), supported shapes, FLOP-counting convention, and whether the rate assumes sparsity—a required pattern of zero inputs. The formats deserve particular attention: fewer input bits reduce data movement, but also change which values survive conversion.

Choose a format for range and precision

A binary floating-point value combines a sign, a significand containing its significant binary digits, and a power-of-two exponent that scales it. The encoding stores exponent and fraction fields; for normal values, a leading significand bit is implied. More exponent bits usually widen the range of magnitudes. More fraction bits put representable values closer together, improving precision. FP32, FP16 and FP8 use 32, 16 and 8 bits respectively; BF16, short for bfloat16, also uses 16 bits. Their field sizes explain the differences below.

FP16 and BF16 both use 2 bytes. Near 1, FP16's gap is 8 times smaller (2⁻¹⁰ versus 2⁻⁷), but BF16 can approximate 70000 while FP16 overflows. Choosing between them requires both the input range and the error the application can tolerate. The FP8 names identify different field sizes: E4M3 has 4 exponent bits and 3 fraction bits; E5M2 has 5 and 2. E4M3FN is a finite-number variant with no infinity encoding, which affects its overflow behaviour.

A value slightly above the largest finite number can still round to that number: 65505 rounds to FP16's 65504. For FP8 values that round beyond the finite range, the explorer offers two conversion modes. Non-saturating conversion produces NaN (not a number) for E4M3FN and infinity for E5M2. Saturating conversion limits the magnitude to 448 or 57,344, preserving the sign. Compare 460 and 470 in E4M3FN: the first rounds to 448, while the second overflows without saturation. FP32, FP16 and BF16 also encode positive and negative infinity. The explorer accepts finite inputs and omits signed-zero distinctions and some hardware conversion details.

Absolute error measures the difference |stored − x|. For nonzero x, relative error divides it by |x|. The same absolute error is therefore a larger proportion of a smaller input. At x = 0 that ratio is undefined, so the explorer reports an exact stored zero instead. A nonzero input rounded to zero has 100% relative error even if the absolute difference is tiny. Subnormal values extend the range below the smallest normal value before conversion reaches zero.

Represent small values with subnormals

A normal FP16 value has an implicit leading 1 before its 10 stored fraction bits, giving 11 significant bits. The smallest positive normal value is 2⁻¹⁴ (about 6.1 × 10⁻⁵). Below it, subnormal values use a leading 0 and fixed spacing of 2⁻²⁴. Their magnitudes are integer multiples of that step, down to the smallest positive FP16 value, 2⁻²⁴ (about 6 × 10⁻⁸). As the multiples get smaller, fewer significant bits remain.

The encoding permits this gradual approach to zero. Arithmetic instructions can handle subnormals differently: some preserve them, while others replace tiny inputs or results with zero, a behaviour called flush to zero. Processing speed can also differ. Check the instruction, format and execution settings when small values matter; the conversion explorer does not model an entire arithmetic pipeline.

Choose how to accumulate the products

Adding products to an accumulator can introduce rounding beyond the loss from converting inputs. Matrix units often combine narrow inputs with a wider accumulator to retain more of the partial sum. This reduces further loss; it cannot recover an input already converted to zero. The stored accumulator type does not specify every internal rounding step, so check the arithmetic rules of the matrix instruction you use.

The effect of rounding is visible even in ordinary FP32 addition. From 2²⁴ upward in that exponent range, neighbouring values are two apart. The exact sum 2²⁴ + 1 lies halfway between 2²⁴ and 2²⁴ + 2. Round-to-nearest, ties-to-even chooses 2²⁴ because its last significand bit is zero; the other candidate's last bit is one. Both values happen to be even integers, so integer parity cannot decide this tie.

With each addition or subtraction rounded to FP32, (2²⁴ + 1) − 2²⁴ evaluates to 0, while 2²⁴ + (1 − 2²⁴) evaluates to 1. A parallel reduction can group the same additions differently, producing different intermediate rounding and final bits. Test both the result's accuracy and its reproducibility when changing the reduction or hardware.

Choose a scale for integer quantization

Quantization maps inputs to a limited set of representable values. One method stores a small integer q and reconstructs the value by multiplying q by a shared positive scale, s. In this b-bit integer quantizer, the codes run symmetrically from −Q to Q, where Q = 2^b−1 − 1 (7 for 4 bits). Divide each input by s, round to an integer, then clip any code outside the range to the nearest endpoint:

At 4 bits and s = 0.5, neighbouring levels are 0.5 apart and the largest reconstruction is 7 × 0.5 = 3.5. The input 6.2 clips to 3.5, while 0.3 rounds to 0.5. With s = 1, 6.2 rounds to 6 without clipping, but 0.3 rounds to zero. A larger scale accommodates large values by making the spacing coarser. The codes −7…7 use 15 of the 16 possible 4-bit patterns; −8 is unused to keep this range symmetric. Unlike this integer quantizer, floating-point formats have spacing that changes with magnitude.

A tensor is an array of numbers, such as a vector or matrix. If all its values share one scale, an outlier can require spacing too coarse for the smaller values. Per-channel scaling gives each feature group, such as a matrix column, its own scale; per-block scaling does the same for small groups of elements. Both require storing and applying more scales. Some narrow floating-point matrix formats also use block scales. Evaluate the final application result with representative inputs: error acceptable for inference, which uses a trained model to make predictions, may impair training, which updates the model, or a solver that repeatedly refines a numerical solution.

To evaluate a matrix operation, check which input values survive conversion, how partial sums are computed, and how the operands fit the supported tiles. Count useful arithmetic and data movement, then measure both runtime and numerical error.

Assign work to different processing engines

The camera pipeline below uses several kinds of processor. A CPU runs general software and makes control decisions. A GPU can apply an operation to many data items in parallel. A neural processing unit (NPU) provides specialized hardware for supported neural-network operations, such as matrix arithmetic. A direct memory access (DMA) engine moves data after software sets up a transfer, saving the CPU from executing a load/store loop for every element. A system with different kinds of processing engine is called heterogeneous.

A system-on-chip (SoC) integrates processing engines, memory interfaces and input/output circuitry on a chip. Its internal fabric is the network of links that carries requests between them. In this example, the engines reach DRAM through shared memory controllers. Each numbered stage produces data for the next; transfers and coordination between stages contribute to the time for the complete pipeline.

Distinguish the ways memory can be shared

The figure shows several engines using memory. To program them correctly, establish where the storage resides, which addresses each engine can use, how cached copies stay consistent and whether data moves between memories. These are separate properties:

For example, CPU and GPU code may use the same pointer value while the allocation's physical pages reside in memory attached to either processor. In CUDA, unified virtual addressing supplies a common naming scheme; Unified Memory manages access to supported allocations and may migrate their pages. The behaviour depends on the platform and its reported capabilities. An input/output memory management unit (IOMMU) translates device addresses to physical memory addresses and enforces access permissions. These mappings let a device reach storage; a separate synchronization mechanism establishes when the data is ready.

Transfer ownership after completion

Consider the camera frame in the figure. The input engine must finish writing it, and those writes must become visible to the NPU, before the NPU starts reading. The input engine must then leave the buffer unchanged until the NPU has finished using it. Software can track these transitions with completion events that record when particular work has finished. Completion of the NPU's reads allows the input buffer to be reused; completion of its output writes allows the GPU to consume that output. These rules assign ownership: who may modify each buffer, and when.

On a non-coherent path, a cache may retain old data after another engine writes memory. The platform's handoff API may therefore need cache maintenance: writing modified cache lines to memory or invalidating old copies, as appropriate for the transfer. A flag alone does not perform that work. On a coherent path, cached copies follow the coherence protocol, but the program must still order the producer and consumer and prevent conflicting simultaneous accesses. Use the event or synchronization operation whose scope covers the handoff; an event for one operation does not establish completion of unrelated queued work.

Try it: overlap work on independent chunks

Suppose each input chunk must be copied to a device, processed and copied back. For that chunk, the three stages depend on one another and must run in order. Different chunks can overlap: chunk 2 can copy in while chunk 1 computes, and chunk 1 can copy out while later chunks progress. This is pipelining. Separate engines, enough buffers and enough bandwidth let it continue through the batch.

The schedule assumes fixed stage times, separate engines, enough buffers and no contention. An asynchronous API allows the CPU to continue without waiting for device completion; returning from the call does not establish that two device stages overlapped. Their stream dependencies, available copy engines and buffer properties determine whether overlap is possible. Pinned host pages remain at stable physical locations during a DMA transfer. The platform may require pinned buffers for asynchronous host/device copies. Transfers through separate copy engines can still slow each other if they share a saturated memory interface.

Account for shared power and heat

CPUs, GPUs and accelerators often share a package's limits on power consumption and heat dissipation. Increasing one engine's activity can leave less power available to another, reducing the clock frequency it can sustain. Dynamic power is roughly proportional to activity × capacitance × voltage² × frequency: switching more circuitry, charging larger electrical loads or running at higher voltage and frequency consumes more power. Leakage adds power consumption even when circuitry is not switching.

Compare energy by measuring joules for the whole completed task, including transfers and conversions as well as arithmetic. An NPU's arithmetic may consume less energy while the complete route consumes more. Measure a representative sustained workload too: a short test can finish before the system heats enough to lower its clocks. The next lesson examines how the physical path to memory affects access time and bandwidth.

Draw the data path. For each buffer, mark who may write it, who must finish reading it before reuse, and which event makes the next stage safe to begin.

How dies connect inside a package

Lesson 26 described how several engines share a system. Their physical construction determines the connections available to them. A die is a piece of silicon containing fabricated circuits. A chiplet is a die designed to work with other dies in one system. The package protects the dies, provides electrical connections and attaches them to the board. An interposer is a layer carrying dense wiring beneath neighbouring dies; a silicon bridge provides such wiring over a smaller area. Dies can also be stacked vertically through 3D integration.

The views below show compute dies and memory stacks in one package. Many parallel connections give the processor a wide memory interface, while their short length helps reduce the energy needed to transfer data. Access to memory elsewhere may require a longer route with less bandwidth.

Smaller dies can improve manufacturing yield, the fraction produced that meet requirements: each die is less likely to contain a defect. Chiplets also let designers reuse components or manufacture different functions with different processes. Connecting and assembling these dies adds cost, and accesses crossing between dies may have different performance. The number of dies alone therefore does not tell software how many processors or memory domains it can use. Universal Chiplet Interconnect Express (UCIe) standardizes a die-to-die connection; compatible interfaces still need compatible functions and system integration.

Try it: where HBM's bandwidth comes from

High Bandwidth Memory (HBM) stacks DRAM dies to fit memory into a small footprint and connects to the processor through a wide interface. An interposer can carry many parallel wires in the available space. Each data pin transfers bits at a specified rate, so increasing the pin count raises the total number of bits that can cross the interface per second. The calculator multiplies that count by the rate per pin and converts bits to bytes:

Reaching the four-stack ceiling requires requests to use all four interfaces. The hardware's address mapping determines how requests are distributed; addresses served mostly by one stack cannot use the full aggregate bandwidth. Delivered throughput also depends on controller scheduling, DRAM refresh, the access pattern and competing traffic. Increasing capacity per stack lets a larger allocation fit, but leaves this interface-bandwidth calculation unchanged.

A wide interface helps move many bytes, but one dependent load must still wait for its value. When it misses in all caches, its request passes through the on-chip interconnect, a memory controller and DRAM. If that value supplies the next load's address, the second request cannot begin until the first completes. The next example examines how memory placement changes this waiting time.

Try it: local and remote memory access

In a NUMA (non-uniform memory access) system, the latency and available bandwidth of memory depend on the requester and the data's location. The operating system describes this locality through nodes, which group resources such as CPU cores and memory. In our machine, a core reaches its local node's controller through the local interconnect; reaching the other node's controller requires an additional link. Moving a page between nodes can change this route while leaving its virtual address unchanged.

Our two-node teaching machine runs a thread on A. A load that misses in all caches takes 80 ns from A's DRAM or 140 ns from B's. The thread follows a chain of pointers: each load reveals the next address, allowing only one request at a time. The slider changes the fraction of requests served by B.

Page placement begins when physical storage is assigned. Under Linux's usual local allocation policy, newly allocated heap pages are placed near the CPU that triggers their allocation, if memory is available there. With lazy allocation, the first write commonly triggers this step, giving the policy its “first touch” name. A successful malloc may reserve virtual space or reuse existing pages; it does not guarantee placement near the eventual worker. If initialization on B places an array there, pointer chasing from A may incur the 140 ns remote path on every cache miss.

Having workers initialize the pages they will use can improve locality. Explicit memory policies, insufficient local memory or later page migration can change the placement. Pinning a worker keeps it on specified CPUs but does not itself move its pages. Check the worker's CPU location and the pages' measured nodes to explain the observed traffic; include the cost of migration if you move the data.

Memory attached to one GPU can likewise require a remote path from another GPU. For a frequently accessed allocation, trace requester → caches → fabric → controller → memory, where the fabric is the interconnect carrying requests between components. Identify which links and controllers the requests use, and what traffic competes for those resources. This connects the package drawing to the latency and bandwidth the workload experiences.

Trace an access from its requester to the memory that serves it. Use the route and access pattern to explain its cost, then check which parts of that physical topology software exposes as locality domains.

Divide the work and identify the required exchanges

Adding GPUs gives a job more compute and memory. Whenever one GPU needs a result produced by another, the program must transfer that result and wait before using it. Strong scaling keeps the total work fixed and gives each GPU a smaller share. Weak scaling keeps each GPU's share roughly fixed and grows the total job. A fixed communication time becomes a larger fraction of the iteration as strong scaling reduces each GPU's compute time.

Model training provides a common example. Training adjusts numerical parameters to reduce a loss, a measure of prediction error. For a batch of examples, it computes gradients: derivatives that describe how the loss changes with each parameter. In data-parallel training, GPUs hold copies of the same model parameters, process different examples, and combine their gradients before applying the same update. If every GPU computes a mean over an equally sized local batch, adding those means and dividing by the GPU count gives the combined mean. Model computations operate on tensors, multidimensional arrays; a layer is one stage of the model, and its outputs are called activations. Tensor partitioning divides an operation within a layer among GPUs. Pipeline partitioning assigns different layers to different devices and transfers activations between them. Each partition determines which data must cross device boundaries.

Communication libraries call a participant a rank, an identity numbered within a communication group; here each rank corresponds to one GPU. A collective involves the whole group. Its participants must issue matching calls in the same order, with compatible sizes, data types and operations. A mismatch can produce incorrect results or leave ranks waiting indefinitely.

Choose the collective by its result

The diagram shows each operation's inputs and outputs. A root is a designated rank: broadcast copies its data to every rank, while reduce delivers the combined result to that root. Reduction combines corresponding array elements with an operation such as sum or maximum; it does not concatenate the arrays. All-reduce delivers that reduced array to every rank. Reduce-scatter reduces the array and gives each rank a different completed piece. All-gather collects pieces in rank order and gives the full collection to everyone. These output rules are the operations' semantics. They determine which collective the program needs, while leaving the transfer algorithm to the implementation.

The ring below implements all-reduce with a reduce-scatter phase followed by an all-gather phase. NVIDIA's Collective Communication Library (NCCL) can choose different algorithms based on message size and the physical connections between devices. Those choices affect time and traffic while preserving the required output. A consumer also needs to know when that output is ready: submission queues the collective, and completion means its result can be used.

For an ordinary NCCL call using a blocking communicator, successful host return means the operation was enqueued, not completed; the host may wait for other ranks during submission. Once enqueued, it executes asynchronously, and a later kernel in the same CUDA stream is ordered after it. A consumer in an independent stream needs a dependency, such as an event recorded after the collective. Grouped calls can defer enqueueing until ncclGroupEnd; nonblocking communicators require checking submission status before dependent CUDA calls. Finally, completion on one rank's stream establishes readiness there. It does not establish simultaneous completion on every rank or provide a general barrier for unrelated work.

Try it: a ring all-reduce

Four ranks form a directed ring R0 → R1 → R2 → R3 → R0. Their starting vectors are R0 = [1, 2, 3, 4], R1 = [10, 20, 30, 40], R2 = [100, 200, 300, 400] and R3 = [1000, 2000, 3000, 4000]. Each vector has four 4-byte elements, divided into four one-element chunks. The first desired output is 1 + 10 + 100 + 1000 = 1111; the full result is [1111, 2222, 3333, 4444].

During three reduce-scatter rounds, each rank sends a partial chunk to its next neighbour, which adds it to the corresponding value it already holds. After the third round, each rank owns one chunk containing all four contributions. During three all-gather rounds, ranks forward completed chunks until everyone has all four. Every send uses the state from before that round, so a value cannot traverse several ring edges in a single round.

Each rank began with 4 × 4 = 16 bytes and sent one 4-byte chunk in each of six rounds: 24 bytes sent per rank. Each rank also received 24 bytes. Dividing the vector into chunks lets all four ring edges carry data concurrently, with each round transferring only one chunk per edge.

Derive the communication cost

Replace the four ranks by P ranks, each contributing an M-byte vector. Splitting each vector into P equal chunks makes a send carry M/P bytes. There are P − 1 rounds in each of the two phases, so each rank sends that chunk size 2(P − 1) times:

Here α is a fixed startup time per round and β is the payload bandwidth of each directed ring link. Thus one round takes approximately α + (M/P)/β. The model assumes equal chunks, identical full-duplex links that can send and receive simultaneously, no competing traffic and negligible addition time. Each rank receives while sending, so receive time is not added again. For P = 4 and M = 100 MB, each chunk is 25 MB. Six sends total 150 MB per rank. At β = 25 GB/s and α = 5 μs, that is 6 ms of transfer plus 6 × 5 μs of startup, or 6.03 ms.

Now keep the vector size fixed while adding ranks. Compute per rank may shrink, but each rank sends 2(P−1)/P × M bytes, a quantity that rises toward 2M. The bandwidth term approaches 2M/β from below, and the startup term 2(P−1)α continues to grow. Even this ideal ring's communication time therefore increases with rank count:

The table describes this ideal ring. Actual libraries may divide messages differently or choose another algorithm; shared links and software overhead also affect runtime. The next explorer models strong scaling with a fixed total compute time and fixed gradient size. Speedup is the one-rank time divided by the multi-rank time; efficiency divides that speedup by the rank count.

A logical ring still runs on physical links

The ring formula gives every directed edge its own β. On a physical machine, several edges may share a network card or a switch port whose capacity is lower than the combined traffic feeding it. Those edges compete for that capacity. A node is one machine in the cluster: transfers within it may cross GPU links, PCIe switches or CPU interconnects, while transfers to another node also use network interfaces and switches. The route determines the bandwidth, latency and competing traffic each transfer encounters.

A bandwidth figure needs a stated boundary: per link or aggregate, one direction or both summed, raw line rate or useful payload, and within one node or across a network. A “direct” device-to-device transfer avoids staging through a CPU-memory buffer. It still takes time and requires accessible buffers plus correct synchronization. The route itself does not establish cache coherence or tell a consumer when the result is ready.

Overlap follows the dependency graph

The scaling bars place all communication after all computation. Training software can group parameter gradients into communication buffers called buckets. Once all gradients in an early bucket are ready, its all-reduce can run while later gradients are computed. In this schedule a bucket waits for its producers, and the parameter update waits for the reduced result. The final bucket often finishes too late to hide its communication behind more gradient computation. That waiting remains on the critical path, the dependency chain that determines the iteration's completion time. Communication and compute may also compete for memory bandwidth while they overlap.

For fixed compute and communication durations, their sum gives the serialized time; max(compute, communication) is an optimistic lower bound with perfect overlap. Data dependencies and shared resources determine how close a program can get to that bound. Inspect a timeline through final synchronization and measure the whole iteration. If the transfer algorithm changes the order of floating-point additions, also recheck the numerical result: the rounding effects from Lesson 25 apply to reductions across GPUs too.

Before quoting a bandwidth or claiming a scaling result, draw the collective's data dependencies and the physical routes it uses.

Consider one iteration of training a model. A CPU prepares input in host memory. Four GPUs each receive a chunk and run low-precision matrix operations with a specified accumulator. Each GPU produces a gradient: values used to determine how the model's parameters should change. An all-reduce combines their FP32 gradient elements (4 bytes each), and each GPU returns a compact result to the CPU. The four lessons let us describe the numerical choices, storage, transfers and dependencies along this path.

1. Name the representations. Specify the formats of the inputs, scale metadata, accumulators and communication buffers, including any output conversion. Quantizing the inputs leaves the FP32 gradient buffer unchanged unless that buffer is also converted.
2. Name the physical homes. Identify the CPU node holding the host data and the GPU holding each device allocation. A pointer identifies an address; placement determines the physical route an access takes.
3. Draw every movement. Include host → GPU, HBM → on-chip tiles, GPU → neighbour GPU, and GPU → host. One byte can cross several boundaries, each with its own traffic count.
4. Establish the handoffs. Input produced before transfer; transfer before use; local gradients before the reduction; reduction complete before its result is used; output transfer complete before the CPU reads it.
5. Find the critical path. Trace the longest chain of dependent stages that determines when the result becomes available. Independent chunks can overlap where engines and bandwidth allow. Include initialization, conversion, launches, transfers, collectives and synchronization in the interval you optimize.
6. Validate. Compare the result with a numerical reference using the application's error tolerance, then measure steady-state and first-use behaviour separately.

Transfer exercise: annotate a path and a time budget

Use a fresh hypothetical machine. One CPU node holds a 400 MB input batch; four GPUs each copy a 100 MB share over an independent 25 GB/s host link. Each GPU then computes for 12 ms and produces its own 200 MB FP32 gradient. A ring all-reduce uses β = 50 GB/s per link and α = 10 μs per round. Finally, each GPU copies a 4 MB result to the host over its link, where the CPU reads it. Here MB means 10⁶ bytes.

For the estimate, assume the host links and host memory do not contend, and the four GPUs perform corresponding stages concurrently. A stage waits for the previous stage's completion, so there is no overlap between stages. Compute time and each GPU's returned result size scale with its local input; the gradient size stays 200 MB because it depends on the model. Treat the CPU's final read time as negligible, since no time for it is given.

1. Draw where the input, the gradients and the outputs live, and mark every copy and every completion dependency.
2. Compute each stage's time and the critical-path estimate for one iteration.
3. Name the bottleneck, and propose one measurement that could disprove your prediction.
4. Now use 8 GPUs. Compare strong scaling (same 400 MB batch) with weak scaling (each GPU keeps 100 MB, so the batch doubles). The gradient stays 200 MB because it depends on the model, not the batch.

Further reading and model limits

Every machine, memory stack, link speed and timing on this page is an invented teaching model. To check the real rules:

• Matrix instructions and FP8: NVIDIA's PTX ISA (supported shapes, types and which threads participate), CUDA WMMA, and the FP8 conversion functions (saturating versus non-saturating).
• Unified addressing versus managed memory: the CUDA Programming Guide's memory chapter.
• Collective semantics: NCCL's collective operations, including what reduce leaves on non-root ranks.
• NUMA placement: the Linux kernel's NUMA memory policy documentation.

Based on the Unit 6 outline and worked answers in chatgpt-reference/. The examples use the teaching assumptions stated in each lesson; the sources above document the corresponding hardware and programming rules.