GPU Architecture · Logical Map
The five names you write for storage — and the physical tier hiding behind each. This closes the second map.
Lesson 4 gave you the physical tiers: registers, L1/shared SRAM, L2, DRAM. But in code you never write "L2." You write in terms of memory spaces — CUDA's logical names for where a variable lives. There are five, and the whole skill is knowing which physical tier each one lands on — because sometimes the name lies to you.
Here's the complete set, how you declare each, and — the payoff column — which Lesson-4 tier actually backs it:
| Space | You declare it as | Who sees it | Lives as long as | Physical tier |
|---|---|---|---|---|
| Register | a plain local variable — float x; | one thread | the thread | register file · ~1 cyc |
| Local | a local var that won't fit / an indexed array | one thread | the thread | DRAM ⚠ · ~400+ cyc |
| Shared | __shared__ float t[256]; | one block | the block | L1/shared SRAM · ~30 cyc |
| Global | cudaMalloc + a pointer | all threads + host | until you free it | DRAM · ~400+ cyc |
| Constant | __constant__ float c[64]; | all threads (read-only) | the application | DRAM + constant cache |
Two of these you fully own from earlier: registers (Lesson 2's occupancy story) and shared (the scratchpad slice of L1 you met last exchange). The three that need care are local, global, and constant — and local and constant are where the logical name diverges from the physical truth.
Notice you never name a tier — you write __shared__ or a bare variable, and the compiler + hardware place it. Same pattern as the compute map: you write logical, the machine supplies physical.
The single most misleading name in CUDA. Local memory is per-thread (that's the "local" — local to a thread) — but physically it lives out in DRAM, the slowest tier, cached through L1/L2 like any global access. It is not on-chip.
When does a variable land here? Two cases:
float a[8]; that you index with a runtime value. Registers aren't addressable by index, so the array is forced into local memory.A kernel that looks like it uses only fast per-thread variables can secretly be hammering DRAM because of spills or indexed arrays. "Local" sounds fast; it's the opposite. This is precisely the physical↔logical gap your mission is about — the name is logical, the cost is physical.
Constant memory is a small (64 KB total), read-only region the host fills before launch. It physically lives in DRAM too — but it's served through a dedicated constant cache with one special power: broadcast. When all 32 threads of a warp read the same address, the cache serves it in a single fetch, effectively as fast as a register.
The flip side: if threads in a warp read different constant addresses, those reads serialize. So constant memory is superb for values every thread shares (coefficients, parameters) and poor for per-thread-varying data. It's a tier tuned for one access pattern.
You asked what makes shared memory fast. The answer is banks. Shared memory is split into 32 equal banks (one per lane in a warp), and it can serve one address from each bank simultaneously — 32 values in a single cycle. That parallelism is its speed.
It holds only when the warp's 32 threads hit 32 different banks. Two cases break it:
Top: each thread hits its own bank → one cycle. Bottom: two threads target the same bank (different words) → the hardware serializes them, halving throughput. That's a bank conflict.
The one exception, again, is broadcast: if all threads read the same word in one bank, there's no conflict — one read serves everyone. Laying out shared arrays so a warp strides across banks (not down one) is a core optimization you'll do later. For now: shared memory is fast by construction, and bank conflicts are how you accidentally give that speed back.
This table is the whole course in one frame. On the left, the two logical hierarchies you write; on the right, the physical machine that runs them. You built every row yourself.
| Logical — you write | Physical — the silicon |
|---|---|
| Kernel / Grid | the whole GPU (SMs fed by the GigaThread Engine) |
| Block | one SM |
| Warp (32 threads) | one sub-partition's scheduler + 32 lanes (SIMT) |
| Thread | one CUDA-core lane |
| Register / Local var | register file · (or DRAM, if spilled) |
| Shared memory | on-chip L1/shared SRAM (per SM) |
| Global / Local / Constant | DRAM (via L2 + L1 / constant cache) |
Five lessons ago, GPU vocabulary was a pile of loose terms. Now every one of them sits in a two-column structure: you shape the left column; the hardware supplies the right. Fast GPU code is just picking a left-column layout whose right-column reality is cheap — keeping warps converged, keeping data on-chip, keeping DRAM traffic low. That's the entire game, and you can now reason about it.
From memory — one click locks each answer.
GotchaA thread declares float buf[16]; and indexes it with a runtime value. That array most likely lives physically in —
Registers can't be indexed by a runtime value, so a dynamically-indexed array is forced into "local" memory — which is per-thread in scope but physically DRAM. The misleading name hides a slow tier; this is the classic local-memory trap.
MechanismConstant memory serves a warp fastest — like a register — exactly when the 32 threads —
The constant cache broadcasts: one address read by all 32 threads is served in a single fetch. Divergent addresses serialize, and constant memory is read-only so no thread writes it. Banks are the shared-memory mechanism, not the constant one.
ConceptShared memory serves a warp in a single cycle precisely when the 32 threads' addresses —
Shared memory has 32 banks and serves one address per bank per cycle — full speed when the warp spreads across all 32. Multiple threads hitting the same bank (different words) is a conflict that serializes. Coalescing/consecutiveness is a global-memory concern, not shared.
CUDA C++ Programming Guide — §"Memory Hierarchy" and §"Device Memory Accesses." NVIDIA's definition of the memory spaces and their qualifiers (__shared__, __constant__), plus the bank and coalescing rules. The canonical version of everything above. Read it here.
The core arc is complete — you have the full mental model of how GPU code runs and where data lives. Two natural directions from here, both of which you flagged earlier: (1) Optimization — put the maps to work: real tiling, coalescing, occupancy tuning, and the __syncthreads() internals you asked about. (2) Hands-on — fire up your RTX 5080: run deviceQuery to see every number from these lessons on your actual chip, then write and profile a first kernel. Tell me which you want next.