GPU Architecture · Reference

GPU / CUDA Glossary

The canonical nomenclature for this course. Every lesson uses these exact definitions. Colour tells you which map a term lives on.

Terms marked (preview) are defined ahead of the lesson that teaches them, so the reference is complete. Don't feel you must know them yet.

The core trade

Latency

Time to finish one task. CPUs are built to minimize it (big caches, branch prediction). Lesson 1.

Throughput

Amount of work finished per unit time across the whole pile. The GPU's design goal — it will happily make any single task slower to raise it. Lesson 1.

Latency hiding

Keeping arithmetic units busy while individual operations wait, by having so many threads resident that a ready one can always run in place of a stalled one. The reason a GPU works. Lesson 1.

SIMT — Single Instruction, Multiple Threadsphysical

NVIDIA's execution style: one instruction is issued to a group of 32 threads (a warp) that all execute it together in lockstep, each on its own data and registers. The hardware reality behind the "thousands of threads" model. Lesson 2.

Occupancyphysical

Resident warps on an SM ÷ the SM's maximum warps. Higher occupancy = more warps to hide latency with. Limited by registers/thread, shared memory/block, and the fixed warp-slot count. Lesson 2 + reference card.

Physical — the silicon

GPUphysical

The whole chip: a collection of Streaming Multiprocessors sharing an L2 cache and a large pool of DRAM (HBM/GDDR). Lesson 1–2.

Streaming Multiprocessor (SM)physical

The fundamental compute unit of the GPU — a self-contained processor with its own lanes, warp schedulers, register file, and L1/shared memory. A modern GPU has dozens (RTX 5080: 84; H100: 132). A CUDA block is assigned to exactly one SM. Lesson 2.

Sub-partition (processing block)physical

One of the 4 near-independent engines an SM is divided into. Each has its own warp scheduler, register-file slice (~64 KB), and set of lanes (~32 CUDA cores). Lesson 2.

CUDA core / lanephysical

A single arithmetic unit (one FP32/INT ALU). Not an independent core — it cannot fetch its own instructions; it executes what its warp scheduler dispatches to all 32 lanes at once. One thread runs on one lane. A spec sheet's "CUDA core" count = total lanes (RTX 5080: 84 SMs × 128 = 10,752). Lesson 2.

Warp schedulerphysical

The unit inside each sub-partition that, every cycle, picks a ready resident warp and issues one instruction to its 32 lanes. Latency hiding happens here — it swaps stalled warps for ready ones. One per sub-partition (4 per SM). Lesson 2.

Warp divergencephysical

When threads within one warp take different branch paths. With only one instruction stream per warp, the paths run serially with non-participating lanes masked off and idle — a performance cost. Branches where a whole warp goes one way are free. Lesson 2.

Tensor Corephysical

A specialized unit beside the CUDA-core lanes — one per sub-partition, four per SM — that computes a matrix-multiply-accumulate D = A×B + C on small tiles (canonically 16×16×16) in one operation, with low-precision inputs accumulated in high precision. Warp-level: the whole warp feeds one MMA. Distinct from CUDA cores; the engine behind "AI TFLOPS." Lesson 11.

Logical — the CUDA programming model

Kernellogical

A function you write that runs on the GPU. Launching it (kernel<<<grid, block>>>(...)) starts one grid of threads, all executing that same function body. Lesson 3.

Threadlogical

The finest unit of the CUDA model: one execution of the kernel body, with its own indices and registers. Runs on one physical lane. Lesson 3.

Warplogical ↔ physical

A group of 32 consecutive threads within a block that the hardware executes together in lockstep (SIMT). The bridge between the two maps: a software-visible grouping that is really a hardware execution unit. Lesson 2–3.

Block (a.k.a. CTA — Cooperative Thread Array)logical

A group of threads (≤ 1024) that can cooperate: they share fast on-chip memory and can synchronize. Assigned whole to one SM and stays for its lifetime — never spans SMs. Blocks must be independent of one another. Lesson 3.

Gridlogical

All the blocks (hence all the threads) of a single kernel launch. The grid is spread across all the SMs of the GPU. Can be 1-, 2-, or 3-dimensional. Lesson 3.

threadIdx / blockIdx / blockDim / gridDimlogical

Built-in per-thread variables. A thread's unique global index is blockIdx.x * blockDim.x + threadIdx.x — the single most-used line in CUDA. Lesson 3.

GigaThread Engine (global work distributor)physical

The chip-level hardware scheduler that assigns blocks to SMs as they free up, in waves when there are more blocks than fit. Decides physical placement; you don't. Lesson 2–3.

Memory — physical tiers (Lesson 4)

Registersphysicallogical

Fastest storage (~1 cycle), private to one thread, physically a slice of the SM's register file. Same name on both maps. Lesson 4.

Shared memory ↔ L1logical / physical

Fast on-chip SRAM inside each SM (~128 KB/SM on RTX 5080, ~30-cycle latency). L1 and shared memory are the same physical SRAM; CUDA exposes a programmer-managed slice as shared memory, private to a block. Lesson 4.

L2 cachephysical

A single cache shared by all SMs, between them and DRAM (~64 MB on RTX 5080, ~200-cycle latency). Largely automatic. Lesson 4.

Global memory ↔ HBM/GDDR (DRAM)logical / physical

The GPU's large main memory (16 GB GDDR7, 960 GB/s on RTX 5080; HBM on datacenter parts). CUDA calls the address space global memory. Big but ~400–800 cycles away. Its bandwidth is a hard shared ceiling — the wall the hierarchy fights. Lesson 4.

Performance concepts (Lesson 4 + optimization)

Latency vs. bandwidthphysical

Latency = wait per access (hideable with more warps). Bandwidth = total bytes/sec, a fixed shared ceiling (not hideable). The key distinction of Lesson 4. Lesson 4.

Arithmetic intensity / rooflinephysical

FLOPs done per byte moved from DRAM. Below the break-even ratio (peak-FLOPS ÷ peak-bandwidth ≈ 59 on RTX 5080) a kernel is memory-bound. Most kernels are. Lesson 4.

Tilinglogical

Loading a chunk of global data into shared memory once, then reusing it many times from there — the archetypal way to escape the bandwidth wall. Lesson 4.

Memory coalescingphysical

When a warp's 32 threads access consecutive addresses, the hardware fuses them into one wide DRAM transaction; scattered accesses issue many, wasting bandwidth. Decides how much of the 960 GB/s you actually get. Lesson 4 teaser; optimization later.

Local memorylogical

A CUDA space that is per-thread in scope but physically in DRAM (cached via L1/L2) — not on-chip and not fast, despite the name. Holds register spills and dynamically-indexed local arrays. Lesson 5.

Constant memorylogical

A small (64 KB) read-only space set by the host, physically in DRAM but served by a per-SM constant cache. Fast (register-like) when all 32 threads of a warp read the same address (broadcast); serializes on divergent addresses. Lesson 5.

Shared-memory banksphysical

Shared memory is split into 32 banks; it serves one address per bank per cycle (32 values/cycle). A bank conflict — two threads of a warp hitting the same bank, different words — serializes those accesses. All threads reading one word = broadcast (no conflict). Lesson 5.

Cooperation & synchronization (Lesson 6)

__syncthreads()logical

A block-scoped barrier: every thread waits until all threads in the block reach it, and all shared/global writes before it become visible after it. The primitive tiling is built on. Block-only — there is no cross-block equivalent (blocks are independent). Must be reached by all threads of the block or none, else deadlock. Lesson 6.

Barrierphysical

A synchronization point where a set of threads wait for each other. Implemented in the SM by a per-block arrival counter; a warp waiting there shows up as the stall_barrier stall reason. Lesson 6.

Data racelogical

Unsynchronized concurrent access to the same memory where at least one is a write — e.g. reading a shared slot another warp hasn't written yet. Nondeterministic garbage; the fix is a barrier. Caught by compute-sanitizer --tool racecheck. Lesson 6.

Grid-wide sync = kernel boundarylogical

Since blocks can't wait on each other mid-kernel, the only synchronization across the whole grid is ending the kernel: split phase-1 and phase-2 into separate launches. Lesson 6.

Execution & operations — the launch path (Lesson 7)

nvcc (compiler driver)operational

Not a compiler but an orchestrator: it runs a .cu file down a device path (→ PTX → SASS in a cubin) and a host path (ordinary C++ with the fatbinary embedded), then hands the host part to gcc/MSVC. Lesson 7.

PTXoperational

A virtual, forward-compatible assembly (targets compute_XX). Not executed directly; the driver JIT-compiles it to SASS if a binary meets a GPU it has no SASS for. The portability copy. Lesson 7.

ptxasoperational

The PTX assembler — nvcc's optimizing back-end that turns PTX into SASS for one specific real architecture. Register allocation and most device-code optimization happen here, not in the front end. Lesson 7.

SASSoperational

The real, architecture-specific machine code the SMs execute (targets a real sm_XX; the name is historical, "Streaming ASSembler"). Won't run across a major-generation change. Inspect with cuobjdump -sass. Lesson 7.

Compute capability (sm_XX / compute_XX)operational

A GPU's hardware feature/ISA generation, written as major.minor (RTX 5050 Blackwell = 12.0). compute_XX names the virtual arch PTX targets; sm_XX the real arch SASS targets. It's the number in nvcc -arch=sm_120 — match it to your GPU to run native SASS with no load-time JIT. Lesson 7.

cubin / fatbinaryoperational

A cubin is an ELF file holding a kernel's SASS plus resource metadata (.nv.info: register/shared-mem/param needs). A fatbinary bundles several cubin/PTX images; nvcc embeds it into your host executable (Linux/ELF section .nv_fatbin). Your CPU binary literally contains GPU code. Lesson 7.

JIT compilationoperational

The driver compiling embedded PTX to the current GPU's SASS at load time (and caching it). The fallback that makes an old binary run on a new GPU. Lesson 7.

Runtime API vs Driver APIoperational

The Runtime API (cuda*, libcudart) is the everyday high-level layer with implicit context/module management; it sits on the Driver API (cu*, libcuda, shipped with the driver) which offers finer control. Lesson 7.

CUDA context / primary contextoperational

"A process on the GPU": its own address space, allocations, loaded modules, and streams. The runtime creates a device's primary context lazily on first use (eagerly at cudaSetDevice since CUDA 12.0). The first-launch cost that skews naïve timings. Lesson 7.

Execution configuration / cudaLaunchKerneloperational

The <<<grid, block, shmem, stream>>> chevrons are sugar: nvcc lowers them to a host stub that calls cudaLaunchKernel(func, grid, block, args, shmem, stream). Arguments are copied into a parameter buffer (delivered via constant memory; cap 32,764 bytes on Volta+). Lesson 7.

Asynchronous launchoperational

A kernel launch enqueues work and returns to the host immediately — usually before the kernel starts. Root cause of async timing and the sticky error model (Lesson 8). Lesson 7.

Pushbuffer / GPFIFO / doorbelloperational

How work crosses into the GPU: the user-mode driver writes GPU commands to a pushbuffer, appends a pointer to it in a per-context ring buffer (GPFIFO), and rings a doorbell — a memory-mapped device register (in a PCIe BAR) that the kernel driver mmapped into the process once, so the write is a plain store that needs no syscall per launch. The GPU's host interface reads the ring. Same ring-plus-doorbell pattern as NVMe. (Model from NVIDIA's open driver / hardware manuals; doorbell is Volta+.) Lesson 7.

Data movement, completion & errors (Lesson 8)

DMA engine / copy enginephysical

A dedicated unit on the GPU that moves bytes between host RAM and device memory over PCIe, independent of the SMs and the CPU. How many overlaps it allows is reported by asyncEngineCount (0/1/2). Lesson 8.

Pinned (page-locked) memoryoperational

Host memory the driver locks in place so its physical address is fixed, letting the DMA engine read it directly — skipping the hidden "bounce buffer" staging copy that pageable memory needs. The precondition for a truly asynchronous cudaMemcpyAsync and for overlapping transfers with compute. Allocate with cudaMallocHost/cudaHostAlloc. Lesson 8.

Streamoperational

An ordered, in-order queue of GPU work. Operations within a stream run in enqueue order; operations in different streams may overlap. The unit of both ordering and concurrency. Beware the legacy default stream (no stream arg): it implicitly synchronizes with other streams and silently serializes overlap. Lesson 8.

Eventoperational

A marker enqueued into a stream (cudaEventRecord). Used to wait for a point in a stream (cudaEventSynchronize), express cross-stream dependencies (cudaStreamWaitEvent), and — via cudaEventElapsedTime — time a kernel using GPU-recorded timestamps (the correct way, immune to the async launch gap). Lesson 8.

Synchronization (host↔device)operational

Making the CPU wait for GPU work: cudaStreamSynchronize (one stream), cudaDeviceSynchronize (all), cudaStreamQuery (non-blocking poll). Under the hood the GPU writes a rising semaphore/fence value; the CPU learns of it by spin (busy-poll, ~100% CPU) or block (sleep until a PCIe interrupt), selectable via cudaSetDeviceFlags. Lesson 8.

Unified Memoryoperational

cudaMallocManaged: one pointer valid on host and device, with pages faulting in and migrating on demand (Pascal+). Convenience over explicit copies — but on WSL2, concurrent CPU/GPU access is unsupported. Lesson 8.

Async / sticky error modeloperational

Because launches are async, kernel faults surface later. Launch/config errors return at the launch; execution errors (e.g. illegal address 700) surface on the next sync/call. Execution faults are sticky: they corrupt the whole context, so every later call returns the same error until the process/context is torn down. cudaGetLastError returns and clears; cudaPeekAtLastError doesn't. Lesson 8.

Sharing & multi-tenancy (Lesson 9)

Concurrent kernels / Hyper-Qoperationalphysical

Kernels from one context can overlap only if they're in different non-serializing streams and spare SM resources exist. Hyper-Q (Kepler+) gives the front end up to 32 hardware work queues so streams don't falsely serialize. A kernel that fills the GPU runs alone regardless — overlap mostly appears with small, low-occupancy kernels. Lesson 9.

Time-slicing (context switching)operational

The default way a GPU serves multiple processes: work from different contexts can't run on the compute engine at once, so each context gets a scheduled slice of the whole GPU, one at a time. Sharing is temporal, not spatial. Lesson 9.

MPS (Multi-Process Service)operational

A daemon+server that routes several processes' work through one shared context so their kernels run truly concurrently — filling a GPU a single process would underuse. Trade-off: weaker isolation (a fatal fault can cascade to co-resident clients). Lesson 9.

MIG (Multi-Instance GPU)physical

Hardware partitioning of one GPU into up to 7 isolated instances, each with dedicated SMs, memory, and cache paths — strongest isolation and guaranteed QoS. Datacenter-only (A100/H100/…); not on GeForce. Lesson 9.

CUDA graphoperational

A captured DAG of GPU operations, instantiated once and relaunched with a single CPU call. Amortizes per-launch submission overhead when a workload fires many small kernels in a loop and the CPU becomes the bottleneck. Lesson 9.

Observability & debugging (Lesson 10)

-lineinfo vs -Goperational

Compile flags controlling debug metadata. -lineinfo adds a SASS→source-line map (DWARF .debug_line) with optimizations intact — correct for profiling. -G adds full device debug info and turns optimizations off — needed by cuda-gdb, wrong for profiling. Release builds strip both. Lesson 10.

nvidia-smi / NVMLoperational

The coarse, live view — utilization, memory, power, clocks, throttling, processes — read from driver counters and on-chip sensors. Gotcha: "GPU-Util %" is time-busy, not saturation (one active SM reads 100%). Lesson 10.

Nsight Systems (nsys) / CUPTIoperational

System-wide CPU+GPU timeline — API calls, kernels, memcpies, overlap, gaps. Built on CUPTI (tracing: Activity records with GPU timestamps + Callback API + correlation IDs linking a host launch to its GPU kernel). Low overhead; the "where does wall-clock go" tool. Lesson 10.

Nsight Compute (ncu)operational

Per-kernel deep profiler reading hardware performance counters + PC/warp-state sampling: occupancy, memory-vs-compute (roofline), warp stall reasons (stall_long_scoreboard, stall_barrier…). Uses kernel replay (re-runs the kernel once per counter group) → high overhead; the "why is this kernel slow" tool. Lesson 10.

compute-sanitizeroperational

Correctness via binary instrumentation: memcheck (out-of-bounds/misaligned), racecheck (shared races), initcheck (uninitialized global reads), synccheck (bad __syncthreads()). Finds the bug a lucky run hides; replaces cuda-memcheck. Lesson 6 & 10.

cuda-gdboperational

Source-level device debugger: breakpoints in kernels, focus on a specific block/warp/lane, inspect registers/memory. Needs a -G build. Reads state; doesn't sample counters. Lesson 10.

Tensor Cores & matrix acceleration (Lesson 11)

MMA — matrix multiply-accumulatephysical

The one operation a Tensor Core performs: D = A×B + C where A, B, C, D are small matrix tiles (canonically 16×16×16), not scalars. One warp instruction replaces the inner loop of scalar FMAs a matmul would otherwise run on the lanes. Lesson 11.

Mixed precisionphysical

The Tensor Core's accuracy trick: inputs in a low-precision format (FP16/BF16/TF32/FP8/FP4…) but products accumulated in a wider one (FP32/INT32), so error doesn't compound over a long dot product. Each GPU generation added input types; the accumulator stays wide. Lesson 11.

TF32physical

A Tensor Core input format (8-bit exponent like FP32, 10-bit mantissa like FP16). Lets ordinary FP32 matmul get a Tensor Core speedup with no source changes — the hardware truncates the mantissa going in, accumulates in FP32. Ampere+. Lesson 11.

GEMMlogical

General matrix-matrix multiply (C = αA·B + βC). The canonical dense-linear-algebra primitive, and the shape everything Tensor Cores accelerate reduces to: FC layers, convolutions (im2col / implicit GEMM), attention, and blocked matrix factorizations. Lesson 11.

WMMA (Warp Matrix Multiply-Accumulate)logical

The CUDA C++ API (nvcuda::wmma) for programming Tensor Cores from your own kernel: opaque warp-owned fragment operands and load_matrix_sync / mma_sync / store_matrix_sync. Warp-cooperative — every call ends in _sync. Above it sit libraries (cuBLAS/cuDNN/CUTLASS); below it, inline mma.sync/wgmma PTX. Lesson 11.

Fragmentlogical

A WMMA operand holding one matrix tile, spread across the 32 lanes' registers in a hardware-chosen layout you don't index by hand. You address tiles, not elements. Three roles: matrix_a, matrix_b, and accumulator. Lesson 11.