6.2.7 · D5GPU Architecture
Question bank — Memory hierarchy (global, shared, registers)
Before we start, we build the three ideas every trap on this page leans on — a warp, a shared-memory bank, and a global-memory segment. Read these first; the traps assume them.
Keep warp, bank, and segment in your hands — the traps below turn on them.
True or false — justify
True or false: Registers are the fastest memory, so using more registers per thread always makes a kernel faster.
False. Each register access is fast, but the SM has a fixed register file; more registers per thread means fewer warps fit, so latency hiding collapses and the kernel usually slows down. Speed-per-access and total throughput are different things.
True or false: Shared memory is a kind of cache the hardware fills for you automatically.
False. Shared memory is a scratchpad you fill by hand with explicit loads and
__syncthreads(). The L1 cache is the automatic one; they often share the same physical SRAM but are controlled differently.True or false: Global memory is slow because its bandwidth is low.
False. Global memory has huge bandwidth (hundreds of GB/s) but high latency (400–800 cycles). It's a wide, long pipe: lots of data per second, but each individual byte takes a long time to arrive. Latency and bandwidth are separate axes.
True or false: If two threads read the same shared-memory address, that is a bank conflict.
False. Same-address reads trigger a broadcast: one bank read is shared to all requesting threads, with no serialization. A conflict needs different words landing in the same bank (see the bank definition above).
True or false: Coalescing is about threads accessing consecutive addresses in program order.
False. It is about the 32 threads of a warp together touching few aligned segments (as defined above). The order within the warp doesn't matter — a permutation of addresses inside one 128-byte segment still coalesces into one transaction.
True or false: Padding a shared array to [32][33] wastes memory and should be avoided.
False in context. The extra column shifts each row's start by one word, so the bank mapping no longer sends a whole column to one bank (worked through in figure s01's caption). The tiny memory cost buys a big serialization win.
True or false: A variable declared inside a kernel always lives in registers.
False. The compiler tries to, but arrays it can't index statically, or spills when register pressure is high, land in local memory — which physically lives in slow global memory despite the name.
True or false: Shared memory is shared across the whole GPU.
False. Its scope is one thread block only. Threads in different blocks cannot see each other's shared memory; the only storage all blocks share is global memory (and L2).
Spot the error
Claim: "My kernel uses 32 registers/thread and gets 100% occupancy, so registers aren't limiting me — I can freely add more registers." What's wrong?
You're at the edge. At exactly the budget-fitting point, adding even one register can drop you a whole block's worth of threads (occupancy is a step function, not smooth). Being at 100% means register pressure is maximally tight, not slack.
Claim: "I made every access hit shared memory, so bank conflicts can't hurt me anymore." What's wrong?
Bank conflicts live inside shared memory. Moving data to shared memory removes global-latency problems but can create a brand-new serialization problem if your stride maps many threads to one bank (per the formula above).
Code: out[col*N + row] = in[row*N + col]; for a transpose is called "coalesced because the read in[row*N + col] is consecutive." Spot the flaw.
The read coalesces, but the write
out[col*N + row] jumps by N between neighboring threads → one segment per thread → uncoalesced. Judging coalescing on only one of the two accesses is the classic mistake.Claim: "Occupancy is 75%, so I'm wasting 25% of the GPU — I must raise it." Spot the error.
Occupancy is a means, not the goal. If 75% already provides enough warps to hide all memory latency, pushing higher does nothing — and forcing it (by cutting registers) can spill variables and make things worse.
Code: __shared__ float As[16][16]; then As[ty][tx] = A[...] — a student says "this loads without needing synchronization." Spot the error.
After the cooperative load you must call
__syncthreads() before any thread reads tiles other threads wrote; otherwise a fast thread reads stale/garbage values. Missing barriers are silent, intermittent bugs.Claim: "Strided global access is only 2× slower than coalesced." Spot the error.
With a large stride each of the 32 threads can land in its own segment → up to 32 transactions instead of 1, i.e. up to 32× the memory traffic. The penalty scales with how badly the addresses scatter across segments, not a fixed 2×.
Claim: "L1 cache and shared memory are completely separate hardware, so using one never affects the other." Spot the error.
On many architectures they are carved from the same SRAM pool. Requesting more shared memory can shrink the L1 cache (and vice versa), so they trade off against each other.
Why questions
Why does giving each thread more registers reduce performance even though registers are the fastest storage?
Because the register file is a fixed pot split across all resident threads. More registers each → fewer resident warps → fewer standby warps to swap in during a stall → latency hiding fails. The bottleneck moves from access-speed to warp-availability.
Why do GPUs hide latency by adding threads instead of adding caches like CPUs do?
GPUs bet on massive parallelism: with thousands of threads, whenever one warp waits on memory, another has work ready, keeping the arithmetic units busy. This "hide, don't shorten" strategy needs many warps, which is why occupancy matters so much.
Why is the shared-memory bank width usually 4 bytes and the bank count 32?
Because a warp is 32 threads and a natural word is 4 bytes: if 32 threads each grab one consecutive 4-byte word, they hit 32 different banks → zero conflict, one cycle. The hardware is tuned to make the common access pattern conflict-free (though other generations may pick different widths).
Why does memory coalescing exist as a concept only for global memory and not registers?
Registers are private per thread — there's no shared transaction to combine. Global memory is fetched in fixed aligned segments serving a whole warp, so how the warp's addresses group into segments determines the transaction count.
Why can two kernels with identical arithmetic have wildly different runtimes?
Because runtime is usually dominated by memory behaviour — coalescing, bank conflicts, occupancy, reuse — not the FLOPs. The roofline view: most kernels are memory-bound, so data movement, not math, sets the clock.
Why does tiling in shared memory speed up matrix multiply even though the total math is unchanged?
Tiling raises arithmetic intensity: each value loaded from global memory once is reused many times from fast shared memory, so global traffic drops sharply while the compute stays the same. Fewer trips to the slow pantry.
Edge cases
What happens to occupancy if a kernel uses zero shared memory and very few registers?
Then registers/shared memory don't limit you — occupancy is capped by the hardware ceilings instead (max warps or max blocks per SM). You can hit 100% but no higher; there's always some binding constraint.
Edge case: all 32 threads in a warp read the exact same global address. Coalesced or not?
Fully coalesced — one segment, one transaction, broadcast to all threads. Same-address is the best case for global memory, just as it is (as broadcast) for shared memory.
Edge case: a warp accesses shared memory with all 32 threads hitting bank 0. What is the ==serialization factor ==, and what is its value here?
is the largest number of threads landing on any single bank — formally . Here all 32 pile on bank 0, so and effective latency is multiplied by 32 — the worst possible shared-memory pattern.
Edge case: your kernel spills registers to "local memory." Where does that data physically go, and why is it dangerous?
Into global memory (backed by L1/L2), so a supposedly fast local variable now costs hundreds of cycles per access. It's dangerous because the name hides the location — you think it's on-chip but it isn't.
Edge case: a thread block has only 1 warp (32 threads). Can this hide memory latency well?
Poorly. With a single warp there's nothing to switch to when it stalls, so the SM idles through the whole memory latency. Latency hiding needs many concurrent warps, which is why tiny blocks underperform.
Edge case: matrix stored column-major but you read it row-by-row with consecutive threads. Coalesced?
No — consecutive threads step across rows, which are
N elements apart in column-major layout, so each lands in a different segment. Layout and access pattern must agree; the same code coalesces on row-major and scatters on column-major.Recall Fast self-test
One idea per line — cover the right side, answer, reveal. Registers fast per-access but limited pot ::: more per thread → fewer warps → worse latency hiding. Broadcast vs bank conflict ::: same address broadcasts (free); different words same bank serialize. bank(a) formula ::: floor(a/4) mod 32 — word index wrapped around 32 banks. Serialization factor S ::: max threads hitting one bank; effective latency times S. Coalescing judged on ::: the whole warp's segment footprint, both reads AND writes. Occupancy is a ::: means to hide latency, not a goal to maximize blindly. "Local memory" physically lives in ::: global memory (register spills), so it's slow.
Back to the parent: 6.2.7 Memory hierarchy.