Gates → GPUs · Notes
Stanford CS149 · Lecture 3 companion

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.
Two ideas connect the lecture. For a calculation that uses each value only once, memory bandwidth can set the pace even when arithmetic hardware is mostly idle. To use that hardware well, we also need to understand both the programming abstraction—what the code promises—and the implementation that maps its work onto cores and SIMD lanes.
Part 1 · Latency and bandwidth

Can memory keep up?

A thought experiment, a highway, some laundry, and a timeline that shows exactly when memory becomes the limit.

Section 1

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?

Many arithmetic units share one route to memory. The NVIDIA V100 has 80 streaming multiprocessors (SMs), each with 64 units for 32-bit floating-point arithmetic: 5,120 in all. They share a 6 MB L2 cache and 16 GB of high-bandwidth memory (HBM2) with a peak transfer rate of about 900 GB/s through a 4,096-bit interface. Section 5 compares the data those units could consume with what memory can deliver.

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.

Section 2

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.

Four ways to move more cars. Drive faster (200 km/h): latency halves to 15 minutes and throughput doubles to 4 cars/hour. Add lanes (four lanes): latency is unchanged, and throughput quadruples to 8 cars/hour. Use the road better: let cars follow each other 1 km apart, and 50 cars are on the road at once. Latency is still 30 minutes, but a car arrives every 36 seconds: 100 cars/hour per lane, 400 cars/hour on four lanes.

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.

Strategy:
Loads:

The completion rate changes. Duplicating every resource finishes two loads every two hours. Pipelining uses one washer, one dryer, and one folder; after startup, a load finishes every 60 minutes, limited by the dryer. Each load needs at least two hours for its three stages, but a later load also waits between washing and drying. The widget's “latency of one load” reports stage time rather than that queueing delay. Its timeline keeps the washer busy while loads remain; the status sentence about an idle washer does not describe this schedule.

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.

Section 3

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.
A wider path can deliver more without shortening a trip. Both paths illustrate a two-second trip; the second has twice the delivery rate. The pictured packets indicate relative width, but the labels undercount the total packets required in flight: at four items per second, two-second latency requires an average of eight; at eight per second, it requires sixteen.

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.

“One operation per clock” describes a rate, not necessarily one operation's latency. In a simplified four-stage instruction pipeline—fetch, decode and register read, execute, write back—one instruction can finish each cycle after the pipeline fills, even though each spends four cycles passing through it. Dependencies or stalls can reduce that rate. Unit 3 builds such a pipeline.
Section 4

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:

1X = load 64 bytes
2Y = add X + Xneeds X
3Z = add X + Yneeds Y

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.

The modeled core has two kinds of work. One path issues a math instruction to the arithmetic unit (ALU), while another may issue a load to the load/store (LD/ST) unit in the same clock. Each hardware thread keeps its own execution context, so the core can work for another thread while one request waits.
Preset:
Memory bandwidth:
Math instructions per load:
Max loads in flight:
Request latency before transfer:

Watch both the arithmetic unit and the memory channel. Each column is one clock. Orange marks a math instruction; red marks a clock with no ready math in this model. The middle row counts outstanding loads, and blue at the bottom marks clocks when the channel transfers a 64-byte block. The reported steady-state utilization is measured over a longer run than the displayed timeline.

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.

LimitFormula (per clock)What raises it
Compute capacity1 math instructionMore ALUs, cores, or SIMD lanes
Memory transfermath per load × bytes per clock ÷ bytes per loadMore bandwidth, or more useful math per block
Requests in flightoutstanding 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.

Section 5

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.

Processor:
Useful operations per element:

Compare an ideal arithmetic demand with memory supply. Needed bandwidth = operations per clock × clocks per second × bytes per operation. At the modeled 1.6 GHz, 5,120 V100 FP32 lanes doing one multiply each clock would need about 98 TB/s if every multiply caused 12 bytes of main-memory traffic. Its nominal bandwidth is 900 GB/s, so at most about 0.9% of that multiply rate is sustainable under these assumptions. The eight-core Xeon preset assumes one eight-wide AVX multiply per core per clock at 3.2 GHz and 76 GB/s of memory bandwidth, giving about 3%. The slider asks what changes when the algorithm performs more useful operations on each loaded element before writing it.

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.

PredictAnswer before opening.
  1. 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.

  2. 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.

  3. 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.

Part 2 · Abstraction vs. implementation

What a program means, and how it runs

Using ISPC to separate a programming model's semantics from the machinery that carries them out.

Section 6

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.

Abstraction (semantics)Which results and ordering does the programming model guarantee? Which choices does it leave open?
Implementation (scheduling)How does a compiler or runtime assign that work to threads, execution units, or vector lanes?
Your goalFirst check that the program is correct under the model's guarantees. Then trace how a likely implementation uses the machine's resources.
The trapA program can appear correct under one execution order while relying on an order the model never promised.

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.

Section 7

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, and numer. Values marked uniform, such as denom and sign, are the same across the gang.
  • The call returns to C++ after all instances in the gang complete.
Fork logically, then join. The C++ caller enters one ISPC call; its gang has programCount logical instances (eight here). C++ resumes after the gang completes. The drawing describes program behavior, not eight operating-system threads.
ISPC keywordMeaning
programCountThe number of logical instances in the gang. It has the same value for every instance (it is uniform).
programIndexThis instance's ID within the gang, 0 to programCount − 1. Different in every instance (a varying value).
uniformA 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.

Section 8

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.

Assignment:
programCount:

Same result, different addresses together. N = 32. Rows are loop steps; columns are program instances. Highlighted array cells are the values requested at that step. In the interleaved version, the eight addresses are contiguous, so one packed load can fetch them; an unaligned AVX example is 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.
“Contiguous work per worker” means different things at different levels. Giving each CPU core a block often lets it stream through nearby addresses. Within one SIMD instruction, neighboring lanes execute together; giving them neighboring elements makes a packed load possible. The useful layout depends on which workers act at the same time and which memory operation the hardware can perform.
Section 9

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.

Section 10

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.

Program:
Iteration order:

Change only the order. x has eight elements; three are negative. The two output rows compare a selected order with index order. For absolute_repeat they match. For shift_negative, some cells have competing writers, highlighted values can differ, and question marks mark unwritten cells. The explorer runs iterations sequentially in a chosen order to illustrate possible results; a real parallel race offers no “last writer” guarantee. It does not predict ISPC's exact schedule.

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
}

Eight private sums, then one result. A possible AVX implementation keeps eight partial sums in a vector register. Each group of eight inputs can be loaded and added to those lanes together; 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:

OperationWhat 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.

Three dependent rounds, seven multiplies overall. Here lane 0 pairs with lane 1, then the partial product in lane 2, then the one in lane 4. The explorer takes a shifted value from d lanes above; an actual shift call must use the offset direction specified by ISPC. Hatched lanes are inactive for that round.
Section 11

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.

Eight logical instances, one CPU vector in this example. On an eight-wide CPU target, the compiler can map an instance to each SIMD lane and use masks when their 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 programCount logical instances with different programIndex values. 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.

Section 12

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 foreach without 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—the partial variable in Section 10 is one example—but code no longer assigns each index with programIndex.
  • Express whole-array operations: in a more restricted style, write y = map(doWork, x) or Z = X + Y and 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.

Who chooses the work assignment? Moving right, more of that choice belongs to the compiler or library. The diagram suggests fewer opportunities for manual mapping errors; it does not promise that every program at a higher level is race-free or efficient.
A programming model is a set of promises, not a wiring diagram. Before relying on an execution order, a lane mapping, or a memory access pattern, ask whether the model guarantees it. Then inspect the implementation to understand performance. That habit will carry through the rest of the course.
Review

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.
Test yourselfThe lecture's questions, plus a few more.
  1. 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.

  2. 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.

  3. 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.

  4. 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.

  5. 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.

  6. Why is shift_negative's output undefined, while absolute_repeat's isn't?
    Show answer

    Assuming separate valid x and y arrays, shift_negative can have two iterations write one y element with no promised order; it can leave other elements unwritten. In absolute_repeat, iteration i alone writes y[2i] and y[2i + 1].

  7. If a language hid programIndex and programCount but kept foreach, could varying values still exist outside the loop?
    Show answer

    Yes. In the sum example, partial is declared before foreach, updated with a different partial sum in each instance, and passed to reduce_add afterward. 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.