Glossary

The old rule that each new, smaller generation of chips ran faster without using more power. It stopped working around 2005.

Precisely: The scaling relationship [Dennard et al., 1974] under which power density stays constant as transistor dimensions shrink; its breakdown is why clock frequencies stopped improving quickly.

See also: Computation accelerator

Hardware built to do one kind of job well, giving up the ability to do everything else well.

Precisely: Specialized hardware trading generality for energy efficiency; Hameed et al. [2010] measured gains up to 500×, roughly 10× from vector execution and the rest from avoiding data movement.

See also: Turing Complete, Dennard Scaling

Able to run any program at all, given enough time and memory — not limited to one task.

Precisely: The property that makes a modern GPU a general-purpose processor rather than a fixed-function accelerator, and the reason GPU computing generalizes beyond graphics.

A GPU on its own card, with its own separate memory, plugged into a computer.

Precisely: A GPU connected to the CPU across a bus such as PCIe, with a distinct DRAM space (device memory) from the CPU's (system memory).

See also: Integrated GPU, System memory / device memory

A GPU built onto the same chip as the CPU, sharing one pool of memory.

Precisely: A CPU+GPU on one die sharing a single DRAM space (e.g. AMD's Bristol Ridge APU), which removes explicit copying but introduces a cache-coherence problem between CPU and GPU caches.

See also: Discrete GPU

The CPU's memory and the GPU's memory. They are tuned for different things: the CPU's for answering quickly, the GPU's for moving a lot of data at once.

Precisely: CPU-side DRAM (DDR, latency-optimized) vs. GPU-side DRAM (GDDR, throughput-optimized); integrated parts share one space, often LPDDR for low power.

See also: Latency hiding

Hardware and software that move data between CPU and GPU memory for you, instead of making the programmer copy it by hand.

Precisely: NVIDIA's automatic CPU↔GPU transfer support (from Pascal onward), built on virtual memory support on both sides [Gelado et al., 2010].

The piece of code the CPU hands to the GPU to run.

Precisely: The program a GPU computing application designates for GPU execution, launched with a specified thread count and input-data location. Detailed in Chapter 2.

See also: Driver, SIMT (Single-Instruction, Multiple-Thread)

CPU-side software that packages up a GPU job and tells the GPU to start.

Precisely: Translates the kernel, thread count and data location into a form the GPU reads, places it where the GPU is configured to look, and signals the hardware.

One instruction is carried out by many threads at the same time, each on its own data.

Precisely: The execution model each GPU core runs for a launched kernel. The mechanism — including what happens when threads disagree about a branch — is the subject of Chapter 3.

See also: GPU core, Kernel

One of the many independent engines inside a GPU. Each runs about a thousand threads.

Precisely: NVIDIA calls it a *streaming multiprocessor*, AMD a *compute unit*. Contains first-level instruction and data caches, scratchpad memory and barrier hardware.

See also: Scratchpad memory, Barrier, SIMT (Single-Instruction, Multiple-Thread)

A small, fast memory inside a core that its threads use to pass data to each other.

Precisely: Software-managed on-core storage enabling inter-thread communication within a core; introduced commercially with the NVIDIA GeForce 8 Series to limit off-chip bandwidth.

See also: GPU core, Barrier

A meeting point: every thread waits there until all the others arrive.

Precisely: Fast synchronization operation among the threads running on a single core.

The job of a small nearby cache — catch most requests so fewer have to travel to distant memory.

Precisely: The role of first-level caches in a GPU core: reduce traffic sent to lower levels of the memory system.

See also: Latency hiding

When one group of threads has to wait for memory, the core runs a different group instead, so no time is wasted.

Precisely: Using a large resident thread count to cover the delay of accesses that miss in first-level caches; the reason GPUs tolerate frequent misses rather than avoiding them.

See also: Performance valley, Bandwidth filter

One slice of the GPU's memory system: a channel to memory plus its share of the big shared cache.

Precisely: A memory channel paired with a portion of last-level cache; multiple partitions provide the memory-level parallelism that sustains throughput. Detailed in Chapter 4.

See also: Interconnection network

The on-chip wiring that lets every core reach every memory partition.

Precisely: An on-chip network, often a crossbar, connecting GPU cores to memory partitions. (The Intel Xeon Phi instead distributes last-level cache with the cores.)

See also: Memory partition

Adding threads does not always help. There is a middle range where you have too many threads for the cache to hold their data, but too few to keep the machine busy while they wait — and performance actually drops.

Precisely: The dip in the Guz et al. [2009] analytical model (Fig. 1.3) between the multicore (MC) region and the multithreaded (MT) region, occurring when the thread count exceeds cache capacity but is insufficient to cover off-chip latency.

See also: Latency hiding, MC region / MT region

The two ends of the performance curve: the left where CPUs live (few threads, big cache) and the right where GPUs live (many threads, hiding waits).

Precisely: The multicore and multithreaded regions of Fig. 1.3, separated by the performance valley.

See also: Performance valley

Hardware that performs one built-in sequence of steps. You can feed it data, but you cannot change what it does.

Precisely: A pipeline whose stages are hardwired rather than programmable. Early 3D graphics processors such as the NVIDIA GeForce 256 were largely of this kind.

See also: Shader, Computation accelerator

A small program the graphics chip runs on your behalf — originally to decide the position or colour of things being drawn.

Precisely: A programmable stage of the graphics pipeline. Vertex and pixel shaders were introduced commercially in the NVIDIA GeForce 3 (2001) [Lindholm et al., 2001], and are the capability that made general-purpose GPU computing possible.

See also: Fixed-function hardware, GPGPU

Using a graphics chip to do ordinary computing work, not just draw pictures.

Precisely: General-Purpose computing on Graphics Processing Units. Began as academic work mapping computation onto shader hardware — first by disguising matrices as textures [Krüger and Westermann, 2003], then behind abstractions that hid graphics entirely [Buck et al., 2004].

See also: Shader, Turing Complete

The ability for the GPU to start new threads by itself, instead of waiting to be told by the CPU.

Precisely: Support for launching threads from the GPU itself. Note that the CPU still initiates the original kernel launch — this adds the ability to spawn further work once running, not to begin work unprompted.

See also: Kernel, Driver

Parts of the chip built specifically to make machine learning fast, rather than to be good at everything.

Precisely: Hardware units introduced in NVIDIA's Volta architecture, targeted at machine learning acceleration. A deliberate move back toward specialization, and the same efficiency-versus-flexibility trade described in §1.1.

See also: Computation accelerator

One thread doing one thread's worth of work, on ordinary single values. This is what you actually write code for.

Precisely: In CUDA and OpenCL, individual threads execute instructions whose operands are scalar values (e.g. 32-bit floating point). The programming model presents these as independent; the hardware groups them at runtime.

See also: Warp / wavefront, SIMT (Single-Instruction, Multiple-Thread)

The group the hardware actually runs together, in step. You write threads; the machine handles them in batches.

Precisely: A group of scalar threads executed in lockstep on SIMD hardware. NVIDIA warps are 32 threads; AMD wavefronts are 64. The grouping happens at runtime and is largely invisible in source code.

See also: Thread block (CTA), SIMT (Single-Instruction, Multiple-Thread), Scalar thread

A team of threads that can share a small fast memory and wait for each other. Bigger than a warp, smaller than the whole launch.

Precisely: Cooperative Thread Array. Warps grouped into a larger unit; threads within one CTA can communicate through per-core scratchpad memory and synchronize with hardware barrier instructions.

See also: Warp / wavefront, Grid, Shared memory / LDS

The whole job: every thread block launched by one kernel call.

Precisely: The top of the thread hierarchy — a grid of thread blocks, each made of warps. Grids and blocks have x, y and z dimensions, set by the kernel configuration syntax.

See also: Thread block (CTA), Kernel configuration syntax

The `<<< >>>` part of a CUDA launch, which says how many thread blocks to run and how many threads go in each.

Precisely: `saxpy<<<nblocks, 256>>>(...)` launches one grid of `nblocks` thread blocks of 256 threads each. Arguments are distributed to every running thread instance.

See also: Grid, threadIdx / blockIdx / blockDim

How a thread finds out who it is: its number inside its block, its block's number, and how big a block is.

Precisely: Built-in CUDA variables giving a thread's coordinates within its thread block, its block's coordinates within the grid, and the block's extent. Combined as `i = blockIdx.x*blockDim.x + threadIdx.x` to assign each thread a unique data element.

See also: Grid, Scalar thread

A tiny standard calculation — multiply a list of numbers by a constant and add a second list. Used here because it is small enough to follow completely.

Precisely: Single-precision A times X Plus Y, from the BLAS library [Lawson et al., 1979]. The chapter's worked example, shown in C, CUDA, PTX, SASS (Fermi and Pascal), OpenCL and GCN microcode.

A small, fast scratch space inside one core that a team of threads can use to pass data around, instead of going out to main memory.

Precisely: Per-compute-core software-managed scratchpad, 16–64 KB per SM, exposed as a distinct memory space (CUDA `__shared__`). NVIDIA calls it shared memory; AMD's GCN calls it the local data store (LDS). Acts as a software-controlled cache and pays off when reuse is frequent and predictable.

See also: Scratchpad memory, Thread block (CTA), Global data store (GDS)

An AMD-only scratch space shared by every core on the chip, not just one.

Precisely: A scratchpad memory in AMD's GCN architecture shared by all cores. NVIDIA GPUs have no equivalent.

See also: Shared memory / LDS

Memory every thread can reach — the only way threads in different teams can talk to each other, and the slow way.

Precisely: The address space accessible to all threads. Threads in different CTAs must communicate through it; access is more expensive than shared memory in both time and energy.

See also: Shared memory / LDS

Letting GPU code start more GPU work itself, so the machine does not sit half-idle when some threads have much more to do than others.

Precisely: Introduced with the Kepler generation. Motivated by load imbalance in data-intensive irregular applications leaving GPU hardware underutilized — the same motivation as Dynamic Warp Formation (§3.4).

See also: Dynamic parallelism

A halfway language. Your CUDA is compiled into it, but no GPU actually runs it — it gets compiled again before it executes.

Precisely: Parallel Thread Execution ISA, NVIDIA's fully documented *virtual* instruction set, introduced with CUDA in 2007. RISC-like, resembling a compiler intermediate representation, notably in its limitless set of virtual registers. Its full documentation is what made GPGPU-Sim possible.

See also: SASS, HSAIL

The instructions the GPU really runs. NVIDIA does not fully publish what they mean.

Precisely: "Streaming ASSembler", NVIDIA's hardware ISA. PTX is compiled to SASS by the driver or by `ptxas`. Only partially documented — as of April 2018, opcode names but no operand formats or semantics. This is deliberate: no published hardware ISA means no backwards-compatibility obligation, freeing NVIDIA to redesign the ISA each generation.

See also: PTX, Control instruction

Attaching a condition to an instruction, so each thread carries it out only if the condition holds for that thread.

Precisely: Conditional execution per lane [Allen et al., 1983], used by both PTX (`@%p1 bra`) and SASS (`@P0`, `@!P0`). One of the important similarities between the virtual and hardware ISAs.

See also: exec mask, SIMT (Single-Instruction, Multiple-Thread)

An extra instruction that tells the hardware how long to wait, so the chip does not need separate circuitry to work out what depends on what.

Precisely: Introduced in NVIDIA's Kepler architecture to eliminate explicit dependency checking with a scoreboard. On Maxwell there is roughly one per three regular instructions; each 64-bit control instruction holds three 21-bit groups encoding a stall count, a yield hint flag, and write/read/wait dependency barriers.

See also: SASS, Operand reuse cache

A small store that lets the chip read the same register several times without going back to the big register file each time — saving energy.

Precisely: Apparent from the `.reuse` flags on regular instructions starting with Maxwell. Lets register values be read multiple times per main register-file access, reducing energy and/or improving performance. Related research in §3.6.1.

See also: Control instruction

AMD's GPU architecture family. Unlike NVIDIA, AMD published exactly what its hardware instructions do.

Precisely: Graphics Core Next; Southern Islands was its first generation. AMD released a complete hardware-level ISA specification [AMD, 2012], which helped researchers build lower-level simulators.

See also: Scalar vs. vector instruction, HSAIL

AMD's equivalent of PTX — a halfway language in their compilation flow.

Precisely: A virtual instruction set architecture in AMD's compilation flow, part of the Heterogeneous System Architecture (HSA).

See also: PTX, GCN / Southern Islands

On AMD chips, some instructions work out one answer for the whole group of threads, and others work out a separate answer for each thread.

Precisely: A key difference between GCN and NVIDIA GPUs. Each GCN compute unit contains a scalar unit coupled with four vector units. `v_` vector instructions compute a different 32-bit value per thread in a wavefront; `s_` scalar instructions compute a single 32-bit value shared by all of them. Useful because parts of a SIMT program often compute the same result regardless of thread ID (§3.5).

See also: GCN / Southern Islands, exec mask

A row of on/off switches, one per thread, saying which threads in the group are currently doing anything.

Precisely: An AMD special register predicating execution of individual vector lanes for SIMT execution. Manipulated by scalar instructions to implement control flow. The masking mechanism itself is covered in §3.1.1.

See also: Predication, Scalar vs. vector instruction

An AMD instruction meaning "wait here until enough of my outstanding memory requests have come back."

Precisely: Enables data-dependency resolution for long-latency operations. Each wavefront has three counters — vector memory count, local/global data store count, and register export count — and the compiler or programmer inserts `S_WAITCNT` to wait until outstanding operations fall below a specified threshold.

See also: GCN / Southern Islands, Latency hiding

A stack of notes the hardware keeps, saying which threads are still to run which piece of code, and where they will all meet up again.

Precisely: A stack of predicate masks implementing SIMT execution alongside traditional predication. Each entry holds a reconvergence PC, a next PC, and an active mask. It handles the two cases predication alone handles badly: nested control flow, and skipping a path entirely when no thread takes it.

See also: Active mask, Immediate post-dominator, Branch divergence

One on/off bit per thread, saying which threads in the warp are doing this instruction for real.

Precisely: The per-entry field of the SIMT stack recording which lanes execute. Written as a bit string, e.g. `1110` for a 4-thread warp.

See also: exec mask, SIMT stack

The address where the split-up threads are due to meet back up.

Precisely: SIMT stack field recording where the threads of this entry will rejoin. At a divergent branch, the pre-branch entry's Next PC is rewritten to this point before the new path entries are pushed.

See also: Immediate post-dominator, SIMT stack

When threads in the same group disagree about which way a branch goes — some want the `if`, some the `else`.

Precisely: Threads within a warp taking different targets at a data-dependent branch. The baseline SIMT stack handles it by serializing the targets, which idles lanes and lowers SIMD efficiency.

See also: SIMT stack, Warp compaction

The first place after a branch that every path is guaranteed to reach, whichever way it went.

Precisely: The earliest point in a program where it can be guaranteed at compile time that diverged threads can again execute in lockstep. The default reconvergence point; runtime mechanisms can sometimes reconverge earlier.

See also: Reconvergence PC (RPC), Thread Frontiers

A hang caused by the divergence hardware itself. One thread gets a lock and leaves the loop, which switches it off — so it can never release the lock, and the others spin forever.

Precisely: A circular dependence introduced by stack-based SIMT [ElTantawy and Aamodt, 2016]. The thread that exits a spin loop reaches the reconvergence point and becomes inactive, unable to execute the release; threads still looping remain active at top-of-stack. **The same program would not deadlock on MIMD hardware.**

See also: SIMT stack, Convergence barrier, Independent Thread Scheduling

NVIDIA's Volta-era replacement for the stack, which lets threads that took different paths take turns instead of being strictly serialized.

Precisely: NVIDIA's thread-divergence management approach from Volta, consistent with a convergence-barrier mechanism disclosed in a patent application [Diamos et al., 2015] and similar to academic Multi-Path IPDOM.

See also: Convergence barrier, Warp split

Instead of a stack, each warp keeps a note of which threads still have to arrive at a meeting point before they can go on together.

Precisely: Per-warp state replacing the SIMT stack: a barrier participation mask (which threads participate), barrier state (which have arrived), and per-thread state, rPC and active bit. Managed by `ADD`, `WAIT` and `YIELD` instructions.

See also: Warp split, SIMT deadlock

A group of threads from one warp that are all at the same instruction, after the warp has divided.

Precisely: A subset of a diverged warp sharing a common PC. Unlike a stack-based implementation, a convergence-barrier scheduler is free to switch between warp splits, which is what enables forward progress when some threads hold a lock.

See also: Convergence barrier, Multi-path execution

Give every warp a turn in a fixed order, so they all make roughly equal progress.

Precisely: Warps ordered fixedly and selected in that order. If warps × issue time exceeds memory latency, execution units stay busy. Favours workloads where threads share data at similar points; disfavours those touching disjoint data.

See also: Latency hiding

A holding area for instructions after they are fetched, so the machine can look at them before deciding whether they are ready to run.

Precisely: Storage after instruction-cache access, typically one or more instructions per warp. Enables dependency information to be examined before issue, and helps hide instruction-cache miss latency in combination with MSHRs.

See also: Scoreboard, Instruction replay

A record of which registers are waiting to be written, so an instruction that needs one knows to wait.

Precisely: GPUs use *in-order* scoreboards, the cheapest design. The naive form needs 8192 bits per core (128 registers × 64 warps) and, if every warp could probe it, 256 read ports. Coon et al.'s design instead keeps a few register-identifier entries per warp, checked when an instruction enters the instruction buffer, producing a dependency bit-vector cleared at writeback.

See also: Instruction buffer, Control instruction

The register file is split into separate compartments. Two reads from the same compartment in one cycle cannot both happen.

Precisely: Single-ported SRAM banks accessed in parallel to simulate a many-ported register file, since SRAM area scales with port count. Modern GPU register files are around 256 KB per core.

See also: Bank conflict (register file), Operand collector

Two values an instruction needs happen to live in the same compartment, so it takes two cycles instead of one.

Precisely: Two operands mapping to the same register-file bank in one cycle. With the naive layout of Fig. 3.13, three instructions take six cycles to read their operands while most banks sit idle. **Not the same as a shared-memory bank conflict (§4.1.1).**

See also: Operand collector, Swizzled register layout

A set of waiting rooms, one per instruction, where operands are gathered as they become available — so several instructions can collect at once instead of queueing.

Precisely: Collector units replace the staging registers of the naive design. Each instruction is allocated one on entering register read; multiple units let several instructions overlap operand reads, giving the arbiter more bank-level parallelism. Forms the third scheduling loop.

See also: Register bank, Swizzled register layout, WAR hazard (operand collector)

Give each warp a different starting compartment, so different warps do not all collide on the same one.

Precisely: Equivalent registers of different warps allocated to different banks (r0 of w0 in bank 0, r0 of w1 in bank 1). Does not help conflicts *within* one instruction, but reduces conflicts *between* warps — especially when warps progress evenly under round-robin or two-level scheduling.

See also: Bank conflict (register file)

A later instruction overwrites a register before an earlier one has finished reading it.

Precisely: Possible because the operand collector imposes no order between instructions [Mishkin et al., 2016]. Fixes evaluated: release-on-commit (up to ~2× slowdown), release-on-read (≤10%), and a bloom-filter *bloomboard* (a few percent). Maxwell's read dependency barrier, managed by control instructions, may serve the same purpose.

See also: Operand collector, Control instruction

When an instruction cannot proceed, the GPU puts it back and tries again later, instead of freezing everything behind it.

Precisely: Instructions are held in the instruction buffer and re-issued rather than stalling. GPUs avoid stalling because distributing a stall signal across a huge register file and deep pipeline hurts the critical path, and because stalling one warp would block unrelated warps behind it.

See also: Instruction buffer

Threads doing the same thing are scattered across several warps, each half-empty. Compaction gathers them into fuller warps.

Precisely: A family of techniques (DWF, TBC, Large Warp Microarchitecture) forming new *dynamic* warps from threads of different *static* warps to raise SIMD efficiency under divergence.

See also: Dynamic Warp Formation (DWF), Branch divergence

Build brand-new warps on the fly out of whichever threads happen to want the same instruction.

Precisely: Fung et al. [2007, 2009]. Requires warps to reach the same divergent branch within a short window, making it very sensitive to the scheduling policy. Two pathologies: greedy scheduling can starve threads, and regrouping increases non-coalesced accesses and shared-memory bank conflicts.

See also: Thread Block Compaction (TBC), Warp compaction

Only regroup threads right after a branch splits them, and put them back before the paths rejoin.

Precisely: Fung and Aamodt [2011]. Extends the SIMT stack across all warps in a core, forcing synchronization and compaction at divergent branches and reconvergence points. More robust than DWF, but synchronizing every warp in a core reduces available thread-level parallelism.

See also: Dynamic Warp Formation (DWF)

Instead of running one branch path then the other, run both at once, filling idle lanes.

Precisely: MPM [ElTantaway et al., 2014] replaces the SIMT stack with a warp-split table and a reconvergence table. Dual-Path Execution (DPS) [Rhu and Erez, 2013a] is the simpler two-split version, extending the stack with one extra PC and active mask.

See also: Warp split, Convergence barrier

Instead of running a warp across many lanes at once, run it through one lane over many cycles.

Precisely: [Keckler et al., 2011; Krashinsky, 2011] Time-multiplexes a warp through a single lane, each lane running its own warp. Achieves SIMD efficiency by fetching each instruction once per warp — amortizing control overhead across *time* rather than across lanes in *space*.

A value that is the same for every thread — so storing 32 copies is a waste.

Precisely: A variable with the same constant value for every thread in the kernel. Can be held in a single scalar register and reused by all threads.

See also: Affine variable, Value structure, Scalar vs. vector instruction

A value that goes up by a fixed step from thread to thread — like an array address. Storing the start and the step is enough.

Precisely: A variable whose value is a linear function of thread ID. Representable as a (base, stride) pair, far more compact than the expanded vector — e.g. `&y[idx] = &y[0] + size(int) * threadIdx.x`.

See also: Uniform variable, Value structure

The regularity across threads that makes scalarization possible — most values are either identical everywhere or a simple step apart.

Precisely: [Kim et al., 2013] The uniform and affine regularity exploited by scalarization and affine execution. Detected either by tag propagation (Table 3.1) or by comparing values at write-back.

See also: Uniform variable, Affine variable

Put a tiny, cheap register file in front of the big expensive one, because most values are read once and then never again.

Precisely: Gebhart et al. [2011b] observe up to 70% of values produced are read only once. A small operand register file (ORF) sits above the main register file, cutting accesses to the large, energy-hungry structure.

See also: Operand reuse cache, RegLess

Hand out physical registers only when a thread actually needs one, rather than reserving a full set up front.

Precisely: Tarjan and Skadron [2011] observe many registers go unused while waiting on memory; renaming allows a 50% smaller physical register file or double the threads. Jeon et al. [2015] add metadata instructions encoding when registers can be reclaimed, and note branch divergence must be accounted for when deciding it is safe.

Get rid of the register file altogether and keep only the handful of values actually in use right now.

Precisely: Kloosterman et al. [2017] replace the register file with a compiler-guided operand staging unit. A capacity manager preloads the registers needed for each program region. A 512-entry OSU matched a 2048 KB register file's performance in 25% of the area, cutting GPU energy by 11%.

See also: Hierarchical register file, Affine variable

Lay the code out so the threads that are furthest behind always run next, and divergent paths meet again on their own.

Precisely: Diamos et al. Sorting basic blocks by a priority derived from the control flow graph and always scheduling the threads that are furthest behind gets reconvergence earlier than the immediate post-dominator, without designating a reconvergence point at all — the code layout does the work.

See also: Reconvergence PC (RPC), Multi-path execution

Compare-and-swap, done atomically: look at a location, and only if it still holds the value you expected, replace it — all in one indivisible step.

Precisely: The CUDA atomic compare-and-swap used to build fine-grained locks. §3.1.2's deadlock example spins on it, and Figure 3.5's printed condition inverts the behaviour its own prose describes — reproduced on that page with the discrepancy explained.

See also: SIMT deadlock, Transactional memory (on GPUs)

The warp you get by default: threads with consecutive IDs bundled together once, and never regrouped.

Precisely: In the baseline GPU, threads with consecutive thread IDs are statically fused into warps for lock-step execution. Nothing in the programming model requires this particular grouping, which is exactly why compaction techniques are free to change it.

See also: Warp compaction, Warp / wavefront, Thread block (CTA)

All the threads in a group happen to want data that sits together, so the hardware fetches it in one trip instead of many.

Precisely: When every thread in a warp accesses locations falling within a single L1 data cache block, only one request needs to be sent to lower-level caches. Coalescing happens at request formation in the load/store unit, *before* the cache is consulted.

See also: Uncoalesced access, Affine variable, Bank conflict (shared memory)

The threads want data scattered across memory, so one instruction turns into many separate trips.

Precisely: When threads within a warp access different cache blocks, multiple memory accesses must be generated. The hardware permits this — it is a performance problem, not an error.

See also: Coalesced access

Two threads want different things out of the same compartment of the scratchpad, so they cannot both be served this cycle.

Precisely: Shared memory is SRAM with one bank per lane, each bank having one read and one write port. A conflict arises when more than one thread accesses the same bank on a given cycle *and wants distinct locations within it*. Two threads reading the same address do not conflict. **Not the same as a register-file bank conflict (§3.3)** — different structure, different fix.

See also: Bank conflict (register file), Instruction replay

A list of memory requests the cache is still waiting on, and what to do when each comes back.

Precisely: Plays a role similar to a CPU's miss-status holding registers, recording the block address and the registers to be written — and additionally recording which deferred memory access the instruction pipeline must replay. Supports merging two requests to the same block.

See also: Instruction replay, Fill unit

The part that catches data coming back from memory and works out which waiting request it belongs to.

Precisely: Receives the memory response from the MMU and uses the request's `subid` field to look up its entry in the pending request table, then reschedules the load through the arbiter — guaranteed to hit, because the line is locked in the cache after being placed in the data array.

See also: Pending request table (PRT)

A cache line is split into smaller pieces, each the smallest chunk memory will hand over in one go.

Precisely: A 32-byte subdivision of a cache line, matching the minimum readable size of a recent graphics DRAM chip such as GDDR5. Fermi and Kepler use 128-byte blocks; Maxwell and Pascal divide them into four sectors.

A queue that lets the texture cache look up tags well ahead of actually fetching the data, so a miss has time to be fetched before anyone needs it.

Precisely: Separates the tag array from the data array in an L1 texture cache [Igehy et al., 1998]. The tag array runs ahead of the data array by roughly the round-trip time of a miss, so cache hits and misses experience roughly the same latency.

See also: Texel, Reorder buffer (texture)

One sample from a texture image — the picture being painted onto a surface.

Precisely: A sample within a texture image, located by coordinates computed in the rendering pipeline. Adjacent screen pixels map to adjacent texels, giving the locality texture caches exploit.

Spreading consecutive chunks of memory across different memory partitions so no single one gets all the traffic.

Precisely: How memory traffic is distributed across memory partition units. An NVIDIA patent describes schemes balancing traffic among up to six partitions at 256-byte or 1,024-byte granularity.

See also: Memory partition

The part that does the final blending work when drawing pixels — and, less obviously, the atomic operations CUDA programs use.

Precisely: Raster operation unit within a memory partition. Used for graphics operations such as alpha blending and surface compression, and it also supports the atomic operations of the CUDA programming model, which can implement synchronization across different thread blocks.

DRAM cannot read one value at a time. It pulls a whole row into a holding area first, and reading another row means putting this one back.

Precisely: A small structure into which a page of DRAM bits is read. Reaching it requires precharging the bitlines and an activate operation, during which no data can be read or written — the dead time memory access schedulers exist to minimize.

See also: Row buffer locality, Memory access scheduler

Requests that land close together in time wanting the same DRAM row, so the row only has to be opened once.

Precisely: A sequence of memory requests has row buffer locality if nearby requests access the same DRAM row in the same bank. Requests from one SM tend to have it; interleaving requests from many SMs at a partition destroys it — which motivates the interconnect arbitration policies of §4.4.1.

See also: Row buffer, Memory access scheduler

Hardware that reorders memory requests so DRAM opens each row as few times as possible.

Precisely: Also called the frame buffer (FB). Reorders read and write operations to reduce row-switch overheads. Implemented with two tables: a *read request sorter*, a set-associative structure mapping all reads to one row in one bank to a single pointer, and a *read request store* holding the individual requests.

See also: Row buffer locality

Deliberately skipping the cache, because putting this data in it would do more harm than good.

Precisely: Proposed in several forms — to avoid associativity stalls when all blocks in a set are reserved for outstanding misses [Jia et al., 2014], and to cut queueing latency at the L2 for warps that are not all-hit [Ausavarungnirun et al., 2015].

See also: Intra-warp / cross-warp contention, Protection distance

A warp's own accesses knocking each other out of the cache, versus different warps knocking each other out.

Precisely: Jia et al. [2014] observed that array indexing in many CUDA applications causes a single warp's memory request to miss on addresses mapping to the same cache set under standard modulo set indexing. Caches with low associativity relative to thread count suffer significant conflict misses.

A count that stops a cache line being thrown out until it has had a fair chance of being reused.

Precisely: A CPU bypassing technique adopted by Chen et al. [2014b]: a line is assigned a protection distance on insertion and cannot be evicted until that many accesses have passed. A request finding no unprotected line in a set bypasses the cache instead.

See also: Cache bypassing

Some warps get all their data quickly, some get none quickly, and a warp is only as fast as its slowest access.

Precisely: Ausavarungnirun et al. [2015] characterize each warp as all/mostly hit, all/mostly miss, or balanced. There is little benefit in a warp that is not all-hit, since mostly-hit warps still wait for the slowest access — which motivates bypassing the L2 for every non-all-hit warp to cut queueing delay for the all-hit ones.

See also: Cache bypassing

Building one very large GPU out of several smaller chips on the same package, because a single chip that big cannot be built.

Precisely: Arunkumar et al. [2017] propose extending performance scaling past the slowing of Moore's Law. Combining local caching of remote data, locality-aware CTA scheduling and first-touch page allocation, they come within 10% of an unimplementable monolithic GPU — 45% better than the largest implementable one in the same process.

A warp that is allowed to run and to read its data, but is not allowed to keep any of it in the cache.

Precisely: In Li et al. [2015], a warp holding no token. It executes normally and its requests are served, but it may not allocate lines in the L1 and so cannot evict another warp's data. This is what lets cache footprint be reduced without reducing the number of warps in flight, unlike warp throttling.

See also: Cache prioritization, CCWS (cache-conscious warp scheduling), Cache bypassing

Deciding which warps are allowed to put data in the cache, separately from deciding which warps get to run.

Precisely: Li et al. [2015] assign tokens to warps to control L1 allocation, giving two independent settings: W, the number of warps that may be scheduled, and T (at most W), the number holding tokens. Statically choosing the best W and T gave a 17% improvement over CCWS with static warp limiting; dynPCAL_MTLP and dynPCAL_CCWS learn them at runtime.

See also: Non-polluting warp, CCWS (cache-conscious warp scheduling)

A holding area that puts returning data back into the order it was asked for, so the data array stays a faithful delayed copy of the tag array.

Precisely: In the L1 texture cache of Figure 4.2, memory returns miss data out of order, but the fragment FIFO queues pointers in request order — so responses must be reordered before they reach the data array, or the queued pointers would index the wrong data.

See also: Fragment FIFO, Texel

Every slot in the relevant part of the cache is already reserved for data that has been asked for but has not arrived, so a new miss has nowhere to go and the pipeline stops.

Precisely: Occurs when all blocks in a cache set are reserved to provide space for data that will be supplied by outstanding cache misses. With space allocated on miss detection and a limited number of MSHRs, intra-warp contention can lead to memory pipeline stalls. Jia et al. [2014] bypass the L1 when a miss occurs and no block can be allocated.

See also: Intra-warp / cross-warp contention, Cache bypassing, Protection distance

Handing out work to GPU cores a whole thread block at a time, rather than a thread at a time.

Precisely: A GPU-wide mechanism assigns each threadblock to a SIMT core based on resource availability, in round-robin order in the baseline machine. Warp slots, registers, scratchpad and block slots are all subscribed at threadblock granularity, and blocks are scheduled until at least one resource in each core is exhausted.

See also: Non-preemption of threadblocks, Threadblock-level throttling

Once a block starts on a core it stays there until it finishes; the GPU cannot pause it and put something else on.

Precisely: Because of the large amount of state associated with each threadblock, contemporary GPUs do not preempt their execution. Threads in the block run to completion before their resources can be assigned to another block. Chimera is the research answer.

See also: Chimera, Idempotence (Chimera)

Deliberately giving a core fewer blocks than it could hold, so they stop fighting over memory.

Precisely: Kayiran et al. [2013] assign each core half its maximum threadblocks, then monitor core idle cycles against memory delay cycles. If a core is predominantly waiting on memory, no further blocks are assigned and existing ones may be paused.

See also: CCWS (cache-conscious warp scheduling), Cache prioritization

A runtime that watches how a core is actually behaving and turns three dials: how many blocks run, and how fast the core and the memory clock.

Precisely: Sethia and Mahlke [2014]. Decides from four counters: active warps, warps waiting on memory, warps ready to issue arithmetic, warps ready to issue memory. Runs in energy-saving mode (scale back the under-utilized resource) or performance-enhancing mode (boost the bottleneck), scaling frequencies by plus or minus 15%.

See also: Threadblock-level throttling

Warps are split into a small pool the scheduler actually looks at each cycle and a larger pool it ignores until they are ready.

Precisely: Gebhart et al. [2011c] use it for energy: selecting from a smaller pool shrinks the warp selection logic, and a warp leaves the active pool on a compiler-identified global or texture memory dependency. Narasiman et al. [2011] use it for performance, forming fetch groups that reach long-latency operations at different times.

See also: Fetch group, CCWS (cache-conscious warp scheduling)

A subset of warps scheduled together, so the groups hit their slow memory operations at staggered times and the machine always has someone to run.

Precisely: Narasiman et al. [2011]. Keeping a group together preserves cache and row-buffer locality within it; switching between groups hides the long latencies. Jog et al. form groups from non-consecutive warps instead, to raise DRAM bank-level parallelism.

See also: Two-level scheduling

One warp coming back to data it loaded itself, versus one warp using data another warp brought in.

Precisely: Rogers et al. [2012] show intra-warp locality is the most common form in cache-sensitive workloads, which is why CCWS keeps a separate victim tag array per warp. Compare Jia et al.'s finer three-way taxonomy in section 4.4.2.

See also: CCWS (cache-conscious warp scheduling), Victim tag array

Notice the cache is thrashing, then run fewer warps so each one gets more of it.

Precisely: Rogers et al. [2012] throttle actively scheduled warps using lost-locality feedback from victim tags. Its central comparative result: the scheduler's decision space is far larger than a replacement policy's, and CCWS with LRU improves hit rate more than prior scheduling mechanisms even when those use Belady-optimal replacement.

See also: Victim tag array, Locality scoring system, DAWS (divergence-aware warp scheduling), Cache prioritization

A per-warp list of what was recently thrown out of the cache, so a later miss can reveal that the warp would have hit.

Precisely: On every eviction the victim's tag is written to the evicting warp's private array. On a later miss that warp's array is probed; a match means intra-warp locality was lost, and CCWS assumes the warp might have hit had it held more of the L1.

See also: CCWS (cache-conscious warp scheduling), Intra-warp vs inter-warp locality

A score per warp that grows when that warp keeps losing data it wanted; warps pushed above a line are made to wait.

Precisely: All warps start with an equal score, assuming no throttling. Scores rise as lost locality is detected; a warp pushed over the threshold is prevented from issuing L1 data cache requests. If no further locality is lost the score decays and the throttled warp resumes.

See also: CCWS (cache-conscious warp scheduling), Victim tag array

CCWS, but it predicts how much cache a warp is about to need from the loop it is in, instead of waiting to be hurt.

Precisely: Rogers et al. [2013] build a per-warp cache footprint estimate for warps in loops and throttle pre-emptively. The estimate shrinks as control-flow divergence takes threads out of the loop. They show an unoptimized benchmark coming close to a hand-optimized version under a better scheduler.

See also: CCWS (cache-conscious warp scheduling)

When the memory system is saturated, let one warp own it and let everyone else get on with computing.

Precisely: Sethia et al. [2015]. A memory-aware scheduler (MAS) with equal-priority and memory-access-priority modes, switching when MSHRs and the miss queue fill; in MAP mode only the owner warp issues memory instructions. Paired with cache access re-execution (CAR), a side buffer that allows L1 hits under misses.

See also: Owner warp

The single warp allowed to keep issuing memory requests while the memory system is overloaded.

Precisely: In Mascar's MAP mode, the owner warp issues until it reaches an instruction dependent on a pending memory request. A non-owner warp's re-executed request that misses in the L1 is not relayed to the L2 but reinserted at the tail of the re-execution queue.

See also: Mascar

Work that can be safely started over from the beginning, so throwing it away costs nothing but the time already spent.

Precisely: Park et al. [2015] use a relaxed definition: whether execution has been idempotent from the start of thread execution. Where it holds, a threadblock can be stopped with no context saved at all.

See also: Chimera, Non-preemption of threadblocks

Three different ways to get a block off a core, picked per block according to which is cheapest right now.

Precisely: Park et al. [2015] choose dynamically among a full context save/store, waiting for the block to finish, and stopping it outright where idempotence allows a free restart. An algorithm picks the subset of running blocks to stop with least throughput impact while meeting a user-specified context switch latency target.

See also: Idempotence (Chimera), Non-preemption of threadblocks

On-chip queues of work per lane, refilled from whichever lanes have too much, so no lane sits idle.

Precisely: Kim and Batten [2014] give each lane a small single-ported SRAM holding work IDs, with ISA instructions to push and pull. Redistribution classifies lanes as greedy or needy against a threshold, or sorts them; a global mechanism moves work between cores, and the structure can be virtualized when work exceeds its capacity.

See also: Topological vs data-driven

Launch one thread per element and let most do nothing, or generate work as you discover it.

Precisely: Irregular GPGPU programs often perform best written data-driven, which improves work efficiency and load balancing but suffers memory contention on pushes and pulls without extensive optimization.

See also: Hardware worklist

Deciding which level of a nested loop becomes threads, which becomes warps, and which becomes blocks.

Precisely: Lee et al. [2014a] identify three: a 1D mapping parallelizing the outer loop; a threadblock/thread mapping giving each outer iteration a block; and a warp-based mapping giving each outer iteration a warp. Their premise is that no mapping is universally optimal, so a framework scores the search space per kernel.

See also: Span and split

How many elements one thread handles, and how to cut that down when handling all of them would starve the GPU.

Precisely: span(1) gives each thread one element and exposes the most parallelism; span(all) gives it every element, required when a dimension's size is unknown until launch or the pattern needs synchronization. split(2) halves that, with a combiner kernel launched to aggregate across the splits.

See also: Nested parallel pattern mapping

Letting kernels launched from the device join an existing kernel rather than each paying full launch cost.

Precisely: Wang et al. [2016a] modify the CUDA model so device-launched kernels share hardware queue resources, aggregating with a running kernel of the same code via a linked list of aggregated thread blocks. 1.4x over CDP and 1.2x over tuned non-CDP CUDA.

See also: CUDA Dynamic Parallelism (CDP)

Mark a region as a transaction and let the hardware sort out conflicts, instead of writing fine-grained locks by hand.

Precisely: The GPU case differs in kind, not degree: tens of thousands of small concurrent transactions rather than tens of large ones, tracked at word-level granularity, on a machine where each private cache is shared by hundreds of threads — which drastically reduces the benefit of detecting conflicts through a cache coherence protocol.

See also: Kilo TM, Value-based conflict detection

The first hardware transactional memory designed for a GPU, built to commit hundreds of tiny transactions at once.

Precisely: Fung et al. [2011]. Combines value-based conflict detection with a recency bloom filter to schedule hundreds of non-conflicting transactions for parallel validation and commit, and extends the SIMT stack with R and T entry types to handle the divergence an abort creates.

See also: Value-based conflict detection, Recency bloom filter, Transaction Retry (R) / Transaction Top (T), WarpTM

To check nobody else changed your data, just read it again and compare.

Precisely: Each transaction re-reads existing global memory to validate, eliminating the need for global conflict-detection metadata. It suits the GPU's highly parallel memory subsystem, avoids direct interaction between conflicting transactions, and detects conflicts at the finest granularity — but a naive implementation requires transactions to commit serially.

See also: Kilo TM, Recency bloom filter

A small approximate summary of what everyone is writing, used to find batches of transactions that cannot possibly conflict.

Precisely: Fung et al. [2011] use notions of time and order to compress a large number of small item sets — the write-sets of all committing transactions. A committing transaction queries it for an approximate conflicting set, some of which are false positives. Being approximate keeps it to a few kB, so it stays on-chip.

See also: Value-based conflict detection, Kilo TM

Two new kinds of row on the SIMT stack: one remembering who has to try again, one tracking the attempt in progress.

Precisely: Kilo TM pushes both at tx_begin. R stores the restart information with an initially empty mask and a PC pointing after tx_begin; T tracks the current attempt. At tx_commit a failing thread sets its bit in R; T is popped and a new T pushed from R's mask and PC. When R's mask is empty both pop, and the revealed N entry's PC is advanced past tx_commit.

See also: Kilo TM, SIMT stack

Handle a whole warp's transactions together, so their memory traffic can be merged the way ordinary GPU accesses are.

Precisely: Fung and Aamodt [2013] amortize Kilo TM's control overhead across a warp. It depends on a two-phase parallel intra-warp conflict resolution; once intra-warp conflicts are resolved, validation and commit accesses can be coalesced into wide accesses that suit the GPU memory system.

See also: Kilo TM, Temporal conflict detection (TCD)

Clocks that agree without talking to each other, used to prove a read-only transaction saw nothing stale.

Precisely: Fung and Aamodt [2013] use globally synchronized on-chip timers that run locally without communicating. Timestamps infer the order of a transaction's reads against other transactions' updates, letting conflict-free read-only transactions commit without value-based validation — a large bandwidth saving for data-structure traversals.

See also: WarpTM, Value-based conflict detection

Get permission for a whole neighbourhood of memory at once instead of asking block by block.

Precisely: Power et al. [2013a] apply coarse-grained region coherence to CPU/GPU integration, where directory bandwidth is a significant bottleneck under GPU traffic. Once region permissions are held, most requests skip the directory and can be offloaded to an incoherent direct-access bus.

See also: Utility-based cache partitioning

Give each core the share of the cache it can actually make use of — then correct for the GPU asking far more often.

Precisely: Lee and Kim [2012] extend Qureshi and Patt's scheme to account for the relative ratio of LLC accesses, since GPU cores generate five to ten times more LLC traffic. Without the correction, CPU-centric management biases capacity toward the GPU. They also sample per-core, having some cores bypass the LLC, to detect whether GPU workloads are cache-sensitive at all.

See also: Region coherence, Cache bypassing

Deliberately running fewer GPU threads so the CPU sharing the chip can get at memory.

Precisely: Kayiran et al. [2014] propose a thread-level-parallelism limiting scheme that observes congestion metrics in the shared CPU/GPU memory controller and interconnection network to estimate how many GPU warps should be actively scheduled on each core. A CPU-centric variant watches only controller stalls; a balanced variant also predicts the GPU harm from a moving average of stalling recorded at each multithreading level, and refuses to throttle further when the predicted increase exceeds a user-set threshold k — a proxy for the priority of GPU performance.

See also: Threadblock-level throttling, CCWS (cache-conscious warp scheduling), Utility-based cache partitioning