Latency, bandwidth, and the SPMD abstraction
Lecture 2 showed how many cores, SIMD lanes, and hardware threads can work at once. They still need data. This lecture first asks how quickly memory can deliver it: even thousands of arithmetic units may wait on one shared memory system. It then asks how to express work for those units without writing every machine instruction ourselves. ISPC provides a concrete way to separate a program's meaning from the hardware steps that carry it out.
- You'll be able to
- Distinguish latency from throughput with everyday and hardware examples. Explain when memory bandwidth limits sustained arithmetic throughput, and calculate the limit in a simple model. Explain why enough requests in flight can hide latency without shortening individual requests. Describe ISPC's SPMD model: gangs, program instances,
programIndex,programCount,uniform,foreach. Trace how interleaved and blocked assignments affect SIMD memory instructions, and spot ISPC programs whose output is undefined. - Background
- Lecture 2 (multi-core, SIMD, multithreading) and Lecture 1 (memory, latency, caches). Units 3–5 of these notes go deeper into pipelining, SIMD and GPUs.
- Source
- Follows the structure and examples of Stanford CS149 Parallel Computing, Fall 2023, Lecture 3, "Multi-Core Architecture Part II + Parallel Programming Abstractions" (Kayvon Fatahalian and Kunle Olukotun). The explanations, diagrams and explorers here are original study material, not a copy of the slides.
A thought experiment: multiply two long vectors
Suppose A and B are arrays far larger than the machine's last-level cache, and we stream through them once. To produce C, the processor reads A[0] and B[0], multiplies them, and writes C[0]; then it can do the same for index 1, index 2, and so on. Formally, C[i] = A[i] × B[i]. The elements are independent, so a GPU can assign many of them to different arithmetic units at once. But each multiplication needs two values brought in and one result written out. Can memory supply those transfers fast enough?
We need two different measures to make that comparison. One describes how long a particular request waits; the other describes how much data can arrive over time.
Latency versus throughput, in everyday terms
Latency is the time from starting one task to finishing it. Throughput is the number of tasks completed per unit of time. Adding parallel capacity can raise throughput even when an individual task takes just as long.
The highway
Imagine a 50 km road from San Francisco to Stanford. At 100 km/h, one trip takes 30 minutes: its latency. If the road admits only one car until that car arrives, two cars can finish per hour. Now compare changing the speed, adding lanes, and allowing cars to follow one another.
Driving faster shortens one car's trip. More lanes or closer spacing instead allow more cars in flight at once; each still takes 30 minutes. Processors use the same two forms of overlap: more execution units can handle separate work, and a pipeline can start new work before earlier work finishes.
Try it: laundry
One load of laundry needs 45 minutes in a washer, 60 in a dryer, and 15 to fold: two hours of work along its path. With several loads, try three schedules below. Watch both the time when all loads finish and the interval between successive completions.
The dryer is the bottleneck: finishing it takes longer than finishing either other stage, so the pipeline cannot produce a load more often than once an hour. Similarly, connecting a pipe that carries 100 liters per second to one that carries 50 gives a path that carries at most 50 liters per second. Widening the already faster part cannot raise the end-to-end rate.
The same two ideas for memory
For memory, the “task” is a request for data. A request has a waiting time, and the memory system has a delivery rate:
- Memory latency is the time from requesting data until it is available. We can express it in nanoseconds or in processor clock cycles, but conversion depends on clock frequency: 100 ns is 400 cycles at 4 GHz.
- Memory bandwidth is the amount of data delivered per second, often measured in GB/s. A system can have high bandwidth even though any one request takes a long time.
This is Little's law for a steady flow: average items in flight = delivery rate × average latency. At eight items per second and two seconds per item, sixteen items must be somewhere in the system at once. Multiple outstanding memory requests create that overlap. They do not make any individual request travel faster; they help the memory channel stay busy while earlier requests are still pending.
When memory bandwidth limits the whole processor
Now use a deliberately simple core model to see where the limit moves. Each of many hardware threads requests a 64-byte block, then does two arithmetic steps that depend on the returned data:
The model can issue one math instruction per clock and issue a load alongside it. Its memory channel transfers 8 bytes per clock, so moving one 64-byte block takes eight clocks after a request is ready to transfer. A request also spends some clocks waiting before transfer. Several threads can request different blocks while others wait, up to the selected number of outstanding loads (requests not yet completed). The instructions and block size are schematic; the question is how often data arrives to enable those two math steps.
Start with the lecture preset. Every completed 64-byte load enables two math instructions. The channel needs eight clocks to deliver that block, so even if it transfers continuously, the core receives enough data for only two useful math clocks out of eight: 25% utilization. The core has spare arithmetic capacity, while the channel is already busy all the time. More threads cannot make that channel transfer faster.
Now raise “request latency before transfer” from 4 to 16 clocks. Three outstanding loads are still enough to keep transfers continuous, so the long-run arithmetic rate remains 25%. Beyond that point, three requests no longer cover the delay: gaps appear between transfers and utilization falls. With only one outstanding load, as in the “too few loads in flight” preset, the channel often waits for the next request to become ready. The same latency can therefore be harmless after bandwidth is saturated and costly before it is saturated.
| Limit | Formula (per clock) | What raises it |
|---|---|---|
| Compute capacity | 1 math instruction | More ALUs, cores, or SIMD lanes |
| Memory transfer | math per load × bytes per clock ÷ bytes per load | More bandwidth, or more useful math per block |
| Requests in flight | outstanding loads × math per load ÷ (request latency + transfer time) | More independent outstanding requests |
These are three ceilings on math instructions per clock in the simplified model. The last is Little's-law reasoning: if K requests may be outstanding and each occupies at least L + D clocks before completion, their completion rate cannot exceed K/(L + D). The smallest ceiling limits the achieved rate in this model. Extra ALUs help when arithmetic capacity is the limit; more independent requests help when too little is in flight; once the channel is full, reducing bytes per useful operation or increasing bandwidth is what matters.
Back to the thought experiment: bandwidth is the critical resource
Return to C[i] = A[i] × B[i]. For 32-bit floats, the simplest traffic count is two 4-byte reads plus one 4-byte write: at least 12 bytes for one useful multiply if those values must come from and go to main memory. That is only 1/12 of an operation per byte, a measure called arithmetic intensity. Cache reuse, write policies, and extra transfers can change actual traffic. This simple count lets us ask whether the available memory bandwidth could sustain the chip's arithmetic peak.
The chip with more arithmetic units can thus have a smaller fraction of them busy on this streaming calculation. If both machines sustain their stated bandwidths and move the same 12 bytes per element, the V100's roughly twelvefold bandwidth gives it roughly twelvefold potential element throughput; actual program speed also depends on transfers, launch costs, and other work. The calculation is bandwidth-bound when the rate of useful operations is capped by data delivery rather than arithmetic capacity.
High-bandwidth memory (HBM) places stacked memory close to the GPU on the package to widen the route for data; the V100's nominal HBM2 rate is 900 GB/s. Even that falls far short of the 98 TB/s demanded by the one-multiply-per-element thought experiment. For an algorithm whose memory traffic is the limit, useful strategies include:
- Reuse loaded data. If several useful operations need A[i] and B[i], keep them in registers or cache while doing those operations. Sharing a loaded block between threads can also avoid duplicate requests when the programming model and hardware support it.
- Compare recomputation with another memory trip. Recomputing a simple value may be cheaper than storing it to memory and reloading it later; expensive calculations can reverse that tradeoff.
- Increase useful work per transferred byte. For example, four useful operations on each loaded pair before one result is stored would raise the simple model's intensity from 1/12 to 4/12 operations per byte.
That last strategy leads into the programming question: when several logical workers process an array, which workers touch neighboring elements, and can they reuse what they fetch? Answering it requires distinguishing the work our code describes from the machine instructions that actually move the data.
- In the explorer's lecture preset, a thread does 2 adds per 64-byte load, and memory delivers 8 bytes per clock. What's the best possible ALU utilization? Does doubling the number of threads change it?
Show answer
64 bytes take 8 clocks and feed 2 adds: at most 2 useful clocks in 8, 25%. More threads don't help: they only add more requests to a memory channel that is already busy 100% of the time.
- How much bandwidth would the V100 need to keep all 5,120 ALUs busy on element-wise multiply?
Show answer
5,120 multiplies/clock × 1.6 × 10⁹ clocks/s × 12 bytes/multiply ≈ 98 TB/s, more than 100 times the nominal 900 GB/s available. Under this idealized one-multiply-per-lane and 12-byte traffic model, at most about 0.9% of the arithmetic rate is sustainable.
- Why does raising the latency from 4 to 16 clocks not change the lecture preset's utilization, while the "too few loads in flight" preset is very sensitive to latency?
Show answer
With enough loads outstanding, the memory channel is always busy, so only its transfer rate matters. With too few in flight (one load in that preset, or three loads once the latency exceeds 16 clocks), the channel waits out part of the latency between transfers. Little's law gives an upper bound of (loads in flight × 64 bytes) ÷ (latency + transfer time); this bound falls as latency rises.
The theme: abstraction versus implementation
The vector multiplication in Part 1 has a simple mathematical result, yet performance depends on details such as which nearby values the hardware fetches together. Parallel programming has this recurring split: the abstraction specifies what work means and what behavior is allowed; an implementation decides how to schedule that work on real hardware. The distinction matters for correctness as well as speed. Here the programming model is SPMD (single program, multiple data), and the common CPU implementation uses SIMD vector instructions. The map of parallel machines and programs places them alongside other models.
ISPC, Intel's Implicit SPMD Program Compiler, makes this split visible. We will describe what its program instances mean first, then see how a CPU can run them with SIMD. That order helps explain why two correct assignments of array elements may have very different memory costs.
SPMD in ISPC: gangs of program instances
Use the same array example as Lecture 2. For each input x[i], measured in radians, approximate sin(x[i]) with a finite Taylor series: x − x³/3! + x⁵/5! − … . With a fixed, small number of terms, each element follows the same calculation, and different elements can be computed independently. Here is an illustrative ISPC version:
export void ispc_sinx(uniform int N, uniform int terms, uniform float* x, uniform float* result) { // assume N % programCount == 0 for (uniform int i = 0; i < N; i += programCount) { int idx = i + programIndex; float value = x[idx]; float numer = x[idx] * x[idx] * x[idx]; uniform float denom = 6.0f; // 3! uniform int sign = -1; for (uniform int j = 1; j <= terms; j++) { value += sign * numer / denom; numer *= x[idx] * x[idx]; denom *= (2*j+2) * (2*j+3); sign *= -1; } result[idx] = value; } }
The function starts value at x, then adds −x³/3!, +x⁵/5!, and so on. For example, terms = 2 keeps the first three Taylor terms. This small kernel illustrates parallel structure, not production-quality sine accuracy: error depends on x and the number of terms, and practical libraries also reduce large inputs to a smaller range.
When C++ calls ispc_sinx(N, terms, x, result), the SPMD model describes a group of logical executions of that function:
- The call executes a gang of ISPC program instances. Each follows the same function with its own logical instance ID. “Logical” matters: this does not promise one operating-system thread or core per instance.
- Each instance has its own varying values, such as
idx,value, andnumer. Values markeduniform, such asdenomandsign, are the same across the gang. - The call returns to C++ after all instances in the gang complete.
programCount logical instances (eight here). C++ resumes after the gang completes. The drawing describes program behavior, not eight operating-system threads.| ISPC keyword | Meaning |
|---|---|
| programCount | The number of logical instances in the gang. It has the same value for every instance (it is uniform). |
| programIndex | This instance's ID within the gang, 0 to programCount − 1. Different in every instance (a varying value). |
| uniform | A type qualifier saying a value is shared across the gang, rather than separate for each instance. It affects which assignments and returns are legal, and it can enable scalar code. |
Trace a concrete call with N = 32 and programCount = 8. The uniform loop counter takes values 0, 8, 16, and 24 for the whole gang. At i = 0, instance 0 uses idx = 0 and instance 7 uses idx = 7; at i = 8 they use 8 and 15. Over the call, instance 0 handles 0, 8, 16, 24, while instance 1 handles 1, 9, 17, 25. This is an interleaved assignment. The divisibility assumption in the code keeps every instance's final index below N; a general version would mask or check the leftover elements.
Assigning work to instances: interleaved versus blocked
A second version divides the array into one contiguous block per instance instead:
// assume N % programCount == 0 uniform int count = N / programCount; int start = programIndex * count; for (uniform int i = 0; i < count; i++) { int idx = start + i; float value = x[idx]; ... // same Taylor series as before result[idx] = value; }
Keep N = 32 and eight instances. In the first version, step 0 assigns indices 0 through 7 across the gang. In this blocked version, each instance owns four adjacent elements, so step 0 assigns 0, 4, 8, …, 28. Both versions compute the same result for every index, but the source code chooses different work assignments. On an eight-wide CPU SIMD target, those eight logical instances can occupy the eight lanes of one vector instruction. With a gang wider than the hardware vector, the compiler may use multiple vectors. The addresses requested together—not just the arithmetic—determine the memory instructions it needs.
vmovups (_mm256_loadu_ps). The aligned vmovaps form also requires a suitably aligned address. In the blocked version, addresses are four floats apart, so a compiler may use a gather instruction or several scalar loads. The other two choices illustrate different assignments, not a promise about how ISPC implements foreach.Raising the level of abstraction: foreach
The interleaved loop is efficient for this SIMD mapping, but the programmer had to write i += programCount, form each idx, and assume N divides the gang size. ISPC's foreach lets us state the intended iterations directly:
export void ispc_sinx(uniform int N, uniform int terms, uniform float* x, uniform float* result) { foreach (i = 0 ... N) { float value = x[i]; ... // same Taylor series, using x[i] result[i] = value; } }
The range 0 ... N means indices 0 through N − 1. foreach gives the whole gang those N iterations; the programmer no longer names the instance responsible for each one. For N = 32 and eight instances, the compiler can process nearby indices together in groups of eight, making packed loads possible, and handle a final incomplete group when N is not divisible by eight. The language does not promise an order in which the iterations finish.
The explorer's first two assignments are concrete ways to distribute these 32 indices. “Instance 0 does all” and “dynamic” are thought experiments about other assignments. In particular, the displayed dynamic mode shuffles a fixed group of indices; it does not simulate a real atomic work queue or timing-dependent scheduling. Its “P× slower” message for one active instance compares work counts, not measured runtime. The lesson is which addresses would reach SIMD lanes together, not the compiler's exact choice or an exact speed ratio.
For independent array elements, foreach reads almost like an ordinary loop. The missing promise of iteration order matters when iterations write the same location or share state, as the next section shows.
Reading ISPC programs carefully: semantics traps
Which of these programs is well defined?
Assume x and y are separate arrays with enough space: x has N elements, and y has at least 2N. The two loops below both process every input index, but their writes have different relationships to neighboring iterations.
// absolute_repeat
foreach (i = 0 ... N) {
if (x[i] < 0)
y[2*i] = -x[i];
else
y[2*i] = x[i];
y[2*i+1] = y[2*i];
}
// shift_negative
foreach (i = 0 ... N) {
if (i >= 1 && x[i] < 0)
y[i-1] = x[i];
else
y[i] = x[i];
}
absolute_repeat writes the absolute value of x[i] into y[2i] and y[2i + 1]. Those two output slots belong only to iteration i, so the result is the same whatever order the iterations use. In shift_negative, a negative x[i] may make iteration i write y[i − 1], while iteration i − 1 may write that same slot through its else branch. For example, x[0] = 3 and x[1] = −1 make both iterations write y[0]. foreach promises no ordering that decides which write wins; some y slots are not written at all. The program therefore has undefined output.
Computing a sum: the uniform/varying distinction
The next problem needs cooperation rather than merely independent iterations. Suppose an exported function must return one uniform float to C++: the sum of x[0] through x[N − 1]. In that function body, two tempting attempts fail type checking for different reasons:
// sum_incorrect_1
float sum = 0.0f;
foreach (i = 0 ... N)
sum += x[i];
return sum; // function requires a uniform float
sum is varying by default: every instance keeps a partial sum. That cannot directly satisfy this function's required uniform return type. The varying-to-uniform return is a compile-time error.
// sum_incorrect_2 uniform float sum = 0.0f; foreach (i = 0 ... N) sum += x[i]; return sum;
sum is uniform: one value for the gang. But x[i] is varying, so the update tries to assign different lane results to that one value. ISPC rejects the varying-to-uniform assignment.
Instead, let each instance accumulate its assigned indices in a private partial sum. With N = 32 and an eight-instance gang, each instance adds four values. After the loop, ISPC's reduce_add combines the eight partials into one uniform result:
export uniform float sum_array(uniform int N, uniform float* x) { float partial = 0.0f; // one per instance foreach (i = 0 ... N) partial += x[i]; return reduce_add(partial); // one uniform total }
reduce_add then combines the lanes, possibly entirely in registers. The explorer uses small whole-number inputs, so its total is exact. For general floating-point inputs, this grouping can round differently from a left-to-right scalar sum.reduce_add is a cross-instance operation: it combines the values held by active instances. ISPC also provides operations that find an extreme value or move a value between instances:
| Operation | What it does |
|---|---|
| uniform int32 reduce_add(int32 x) | Sum of the active instances' int32 values; the float overload returns a uniform float |
| uniform int32 reduce_min(int32 a) | Minimum of a over the active instances |
| int32 broadcast(int32 value, uniform int index) | Active instances receive the value held by instance index |
| int32 rotate(int32 value, uniform int offset) | Pass values around the gang by offset positions, wrapping at the ends |
| shift(value, offset) | Like rotate, but values shifted past the end of the gang become 0 |
Try it: a tree reduction across the gang
To see why cross-instance movement helps, multiply eight numbers. A left-to-right chain has seven dependent multiplies. In a tree, four pairs multiply at once, then two pairs of partial products, then one final pair: still seven multiplies in total, but only three dependent rounds (log₂ 8 = 3). Each round obtains a partner's partial product with shift; half as many lanes have useful work in the next round.
shift call must use the offset direction specified by ISPC. Hatched lanes are inactive for that round.How the gang is implemented, and what it doesn't give you
We can now connect the logical model to the CPU doing the work. In the eight-instance example, each instance holds one of the eight partial sums, and reduce_add needs to combine them. A SIMD vector can hold those eight values, update them with packed operations, and reduce them at the end. That is a common implementation of one gang, not eight independent CPU threads.
if branches differ. A single ordinary call runs this gang on one calling thread. A wider gang may require multiple native vectors while remaining one logical gang. Tasks or multiple host callers can put other gangs on other cores.- SPMD model: the gang has
programCountlogical instances with differentprogramIndexvalues. The language defines their values, masks, and cross-instance operations. - CPU SIMD implementation: vector instructions perform corresponding work for several instances together. If instances take different branches, the compiler can mask lanes while executing each path.
- Performance consequence: separate logical instances do not automatically mean separate cores. Address patterns and how many lanes remain active determine how effectively one vector unit is used.
To use multiple CPU cores, the application needs multiple independent gangs. One route is ISPC's task mechanism: launch makes work available to a task runtime, whose workers may run gangs on several cores. Another is to call the exported ISPC function from several C++ threads. Scheduling depends on the runtime and the available work; a task declaration alone does not guarantee a particular core. The useful mental picture is several gangs across cores, with SIMD work inside each gang—the multicore × SIMD combination from Lecture 2.
How high should the abstraction go?
ISPC exposes instance IDs and gang size so we can choose assignments such as interleaved or blocked. That control made the memory behavior in Section 8 visible, but it also makes off-by-one indexing, assumptions about gang size, and conflicting writes possible. A higher-level model can hide some of those decisions. Consider two steps away from manual instance mapping:
- Use
foreachwithout manual instance IDs: the programmer specifies the iteration range and leaves mapping to the compiler. Varying values can still exist before and after the loop—thepartialvariable in Section 10 is one example—but code no longer assigns each index withprogramIndex. - Express whole-array operations: in a more restricted style, write
y = map(doWork, x)orZ = X + Yand let a library choose loops, vector width, and parallel scheduling. NumPy and PyTorch support this style, alongside explicit indexing and other lower-level features.
Hiding a choice can remove one way to make a mistake, such as a hand-written out-of-range index. It can also make a performance choice harder to express. No rung automatically guarantees correct code or fast code: shared-state races and costly data movement remain possible wherever the model permits them. Later lectures revisit these data-parallel styles.
Terms to know, and questions to test yourself
- Latency
- Time for one operation or request to complete.
- Throughput
- Operations completed per unit time.
- Memory bandwidth
- The rate at which memory can supply data, in bytes per second.
- Little's law
- In a steady flow, average items in flight = delivery rate × average latency.
- Arithmetic intensity
- Useful arithmetic operations per byte transferred; reuse can raise it.
- Bandwidth-bound
- Useful operation rate limited by memory's data-delivery rate rather than arithmetic capacity.
- Bottleneck
- The stage with the lowest throughput, which sets the throughput of the whole path.
- Abstraction vs implementation
- What a programming model's operations mean, versus how a particular system carries them out.
- SPMD
- Single program, multiple data: one function run by many logical instances on different data.
- Gang, program instance
- A gang is the group of logical ISPC instances in one call. Each instance has its own varying values; uniform values are shared.
- programCount / programIndex
- The gang's size (uniform) and an instance's ID (varying).
- uniform / varying
- One value shared by the gang / a separate value per instance.
- foreach
- Declares iterations for the whole gang, leaving the assignment to the implementation.
- Packed load vs gather
- A vector load of neighboring addresses versus a vector load from separately chosen addresses. Their costs depend on the hardware and memory pattern.
- Cross-instance operation
- A standard-library call that combines or moves values between instances: reduce_add, broadcast, rotate, shift.
- ISPC task
- One way to submit independent gangs to a runtime that may run them across CPU cores.
- A car trip takes 30 minutes with one car on the road at a time. Adding three lanes raises throughput from 2 to 8 cars/hour. What happened to latency?
Show answer
Nothing: still 30 minutes. More lanes add parallelism, which raises throughput without changing any one car's trip.
- Laundry: wash 45 min, dry 60 min, fold 15 min. Once the pipeline is full, what's its completion rate? If washes start every 45 minutes, what happens to later loads' individual latencies?
Show answer
One load per hour. The dryer takes 60 minutes and sets the completion rate. New washes begin faster than the dryer can accept them, so later loads queue after washing and take more than two hours from their own start to finish.
- A memory system delivers eight items per second, and each request takes two seconds. How many requests must be in flight on average?
Show answer
Little's law gives 8 × 2 = 16 requests in flight on average. The memory sketch's eight icons are schematic; eight literal outstanding requests would not sustain this rate at two-second latency.
- Convince yourself: in steady state, with enough loads in flight, why is core utilization independent of memory latency?
Show answer
Once enough independent requests are outstanding, the channel transfers continuously. Its bandwidth then caps the rate at which operands arrive. Raising latency requires more requests in flight to keep that condition true, but does not lower the already saturated transfer rate.
- For N = 32 and eight instances, why might the blocked sinx version need a gather or separate loads while the interleaved version can use a packed load?
Show answer
At one blocked-loop step, the instances access x[0], x[4], x[8], …, x[28]: addresses four floats apart. At one interleaved step, they access x[0] through x[7], eight neighboring floats that a packed load can fetch.
- Why is shift_negative's output undefined, while absolute_repeat's isn't?
Show answer
Assuming separate valid x and y arrays,
shift_negativecan have two iterations write one y element with no promised order; it can leave other elements unwritten. Inabsolute_repeat, iteration i alone writes y[2i] and y[2i + 1]. - If a language hid
programIndexandprogramCountbut keptforeach, could varying values still exist outside the loop?Show answer
Yes. In the sum example,
partialis declared beforeforeach, updated with a different partial sum in each instance, and passed toreduce_addafterward. Hiding instance IDs removes manual index assignment, not all varying state.
Based on Stanford CS149 (Fall 2023), Lecture 3, "Multi-Core Architecture Part II (latency/bandwidth issues) + Parallel Programming Abstractions", by Kayvon Fatahalian and Kunle Olukotun. The explanations, diagrams and interactive models on this page are original study material. Hardware figures (V100, Xeon E5 v4) and ISPC examples are as presented in the lecture; for ISPC's exact semantics see the ISPC documentation. For the course's own slides and videos, see the CS149 course site.