Memory hierarchy (global, shared, registers)

Overview

GPU memory hierarchy trades off capacity for speed, organizing storage into three main levels: registers (fastest, smallest), shared memory (medium speed, small), and global memory (slowest, largest). Understanding this hierarchy is critical for GPU performance because memory access latency can be 100-1000× longer than computation.

Figure — Memory hierarchy (global, shared, registers)

Core Concepts

1. Registers

Key Properties:

Latency: ~1 clock cycle
Capacity: 32-256 KB per SM (divided among all active threads)
Scope: Private to a single thread
Allocation: Automatic (compiler assigns local variables to registers)

Formula Register Pressure and Occupancy

Derivation from First Principles: An SM has a fixed register budget $R_{total}$ (e.g., 65,536 registers). If each thread uses $r$ registers, the maximum number of concurrent threads is:

$N_{threads} = \left\lfloor \frac{R_{total}}{r} \right\rfloor$

Why this matters: GPUs hide memory latency by switching between threads. If you have fewer threads (because each uses too many registers), the SM can't hide latency → performance drops.

Occupancy (fraction of max threads active):

$\text{Occupancy} = \frac{N_{threads}}{N_{max}}$

Where $N_{max}$ is the hardware limit (e.g., 2048 threads/SM on modern GPUs).

The tradeoff: More registers per thread = fewer concurrent threads = worse latency hiding.

Worked example Register Usage Example

Problem: You launch a kernel with 512 threads/block on an SM with 65,536 registers. Your kernel uses 40 registers/thread. How many blocks can run concurrently?

Solution:

Step 1: Registers needed per block $R_{block} = 512 \times 40 = 20,480 \text{ registers}$

Why this step? We need to know the total demand before checking capacity.

Step 2: Blocks that fit $N_{blocks} = \left\lfloor \frac{65,536}{20,480} \right\rfloor = 3 \text{ blocks}$

Why this step? Each block needs 20,480 registers; we divide total capacity by per-block demand.

Step 3: Active threads $N_{threads} = 3 \times 512 = 1,536 \text{ threads}$ If max is 2048, occupancy = 1536/2048 = 75%

Key insight: If we reduce register usage to 32/thread, we'd fit 4 blocks (2048 threads, 100% occupancy). Even small register savings matter.

2. Shared Memory

Key Properties:

Latency: ~30-100 clock cycles (roughly ~4-13× faster than global memory)
Capacity: 16-96 KB per SM (configurable, often reconfigurable with L1 cache on modern architectures)
Scope: Shared within a thread block
Access pattern: Organized into banks (typically 32 banks, 4-byte width)

Formula Bank Conflicts

Derivation from Memory Banking Architecture:

Shared memory is divided into $B$ banks (typically $B=32$ ). Each bank can service one request per cycle. Address $a$ maps to bank:

$\text{bank}(a) = \left(\frac{a}{w}\right) \bmod B$ Where $w$ is bank width (usually 4 bytes). Why this matters: If multiple threads in a warp access the same bank simultaneously, the accesses serialize → multiplies latency.

Serialization factor:

$S = \max_{i \in [0, B)} (\text{number of threads accessing bank } i)$

Effective latency = base latency × $S$

Bank conflict avoidance: Stride accesses so thread $t$ accesses address $t \times w$ (sequential in 4-byte words).

Worked example Shared Memory Example: Tiling Matrix Multiply

Problem: Multiply two 1024×1024 matrices. Each element requires 1024 multiply-adds. Naive approach: each thread loads from global memory 2048 times.

Solution with shared memory: Step 1: Divide matrices into 16×16 tiles Why? A tile fits in shared memory (~1 KB for float). All threads in the block reuse the same tile data.

Step 2: Load one tile of A and one tile of B into shared memory (coalesced global reads)

__shared__ float As[16][16];
__shared__ float Bs[16][16];
As[ty][tx] = A[row*N + tileStart + tx];  // All threads cooperate

Why this step? 256 threads load 256 elements once, instead of each thread loading the same elements repeatedly.

Step 3: Each thread computes its partial result using the shared tile

for(int k=0; k<16; k++)
    Cvalue += As[ty][k] * Bs[k][tx];

Why this step? Now we reuse each loaded value 16 times from fast shared memory.

Step 4: Repeat for all tiles (1024/16 = 64 tiles)

Speedup calculation:

Naive: 2048 global reads/thread × 1M threads = 2B global reads
Tiled: 2×64 tile loads × 256 elements = 32K global reads per block × 4096 blocks = 131M global reads ~15× reduction in global memory traffic!

3. Global Memory

Key Properties:

Latency: 400-800 clock cycles
Capacity: GBs (e.g., 8-80 GB on modern GPUs)
Scope: Global across entire device
Bandwidth: 200-1000 GB/s (high throughput but high latency)

Formula Memory Coalescing

Derivation from Memory Transaction Model:

Modern GPUs fetch memory in aligned segments of size $S$ (typically 32, 64, or 128 bytes). When a warp (32 threads) requests data:

$N_{transactions} = \text{number of distinct segments touched}$

Best case (coalesced): All 32 threads access consecutive 4-byte addresses

$\text{addresses: } [a, a+4, a+8, ..., a+124]$

Why? All addresses fall within $\lceil 128/128 \rceil = 1$ segment → 1 transaction

Worst case (strided): Thread $t$ accesses $a + t \times 256$

$N_{transactions} = 32 \text{ (each thread in different segment)}$

Performance ratio: 32× difference in memory bandwidth utilization.

The rule: Ensure thread $t$ accesses address $\text{base} + t \times \text{sizeof(element)}$ for consecutive threads.

Worked example Coalescing Example: Transpose

Problem: Transpose a 4096×4096 matrix. Naive approach reads coalesced but writes strided (or vice versa). Naive transpose (bad):

out[col*N + row] = in[row*N + col];  // Read coalesced, write strided

Why bad? Thread $t$ writes to out[col*N + t]. As $col$ stays constant within a warp, all threads write to addresses $N$ apart → uncoalesced. Better solution: Use shared memory tile

Step 1: Coalesced read into shared memory

__shared__ float tile[32][33];  // Note: 33 avoids bank conflicts
tile[ty][tx] = in[row*N + col];

Why 33? If we used [32][32], accessing tile[i][i] for all threads would hit the same bank. The padding shifts banks.

Step 2: Transpose in shared memory (free - just index differently) Step 3: Coalesced write from transposed tile

out[col*N + row] = tile[tx][ty];  // Now coalesced

Result: Both reads and writes are coalesced. ~10× speedup over naive version.

Memory Hierarchy Summary Table

| Level | Latency (cycles) | Size | Scope | Managed By | |-------|------------------|-------|------------| | Registers | 1 | 32-256 KB/SM | Per-thread | Compiler | | Shared Memory | 30-100 | 16-96 KB/SM | Per-block | Programmer | | L1 Cache | 30-100 | 16-128 KB/SM | Per-SM | Hardware | | L2 Cache | ~200 | 0.5-6 MB | Global | Hardware | | Global Memory | 400-800 | GB | Global | Programmer |

Common Mistakes & Misconceptions

Recall Feynman Explanation (Explain to a 12-Year-Old)

Imagine you're doing homework and need to look up information:

Your desk (registers): The paper you're writing on and the pencil in your hand. Super fast to use - you don't even have to reach. But there's only room for one or two sheets.

Your backpack (shared memory): Books and notes you brought to your desk. Takes a second to grab, but way faster than walking across the room. You and your friends working on the same group project can share what's in one backpack.

The bookshelf (global memory): Across the room. Takes 10 seconds to walk there and back. But it holds hundreds of books.

The smart strategy: When you start a new problem, walk to the bookshelf ONCE, grab everything you need, stuff it in your backpack, and pull what you're actively using onto your desk. If you walk to the bookshelf for every single sentence you write, you'll spend all day walking and no time actually working.

That's exactly what GPUs do! They move data from slow global memory → fast shared memory → super-fast registers, and reuse it as much as possible at each level before going back for more.

Connections

GPU Thread Hierarchy - Thread blocks determine shared memory scope
Warp Execution Model - Memory coalescing depends on warp-level access patterns
Cache Architecture - L1/L2 caches complement the explicit hierarchy
Memory Bandwidth Optimization - Techniques to maximize throughput at each level
CUDA Memory Types - Software abstractions (constant, texture memory)
Memory Latency Hiding - Why occupancy matters for performance
Matrix Multiplication Optimization - Classic case study using all three levels
Roofline Model - Analyzing memory vs compute bottlenecks

Flashcards

What are the three main levels of GPU memory hierarchy from fastest to slowest?

Registers (per-thread, ~1 cycle), Shared Memory (per-block, ~30-100 cycles), Global Memory (device-wide, ~400-800 cycles).

What determines how many threads can run concurrently on an SM?

Register usage per thread (occupancy = total_registers / (threads × registers_per_thread)) and other limits like shared memory and thread block size.

What is a bank conflict in shared memory?

When multiple threads in a warp access different addresses that map to the same memory bank, causing accesses to serialize and increasing latency by the number of conflicts.

How does shared memory enable performance optimization?

By allowing threads in a block to cooperatively load data from slow global memory once, then reuse it multiple times from fast shared memory, reducing global memory traffic.

What is memory coalescing?

When consecutive threads in a warp access consecutive memory addresses, allowing the GPU to combine requests into a single wide memory transaction instead of multiple narrow ones.

Why does using many registers per thread hurt performance?

High register usage reduces occupancy (fewer concurrent threads), which reduces the GPU's ability to hide memory latency by switching between threads.

What is the scope of shared memory?

Per thread block - only threads within the same block can access the same shared memory, and data is not visible across different blocks.

Approximately how much faster is shared memory than global memory in latency terms?

Roughly 4-13× faster (~30-100 cycles for shared vs ~400-800 cycles for global), not 100×.

What is the typical size difference between registers and global memory?

Registers: ~100 KB per SM, Global memory: tens of GB - roughly 100,000× larger capacity with 400-800× higher latency.

Concept Map

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Dekho beta, is GPU memory hierarchy ka core idea bahut simple hai — jaise kitchen mein aap cooking karte ho. Aapki pantry (global memory) sabse badi hai lekin wahan tak jaana time leta hai, counter par rakhi cheezein (shared memory) medium fast hain, aur jo cheez aapke haath mein hai (registers) woh instant use ho jaati hai. GPU mein bhi bilkul yahi funda hai — teen levels of memory jahan speed aur capacity ke beech tradeoff hota hai. Registers sabse fast (1 clock cycle) par sabse chhote, shared memory medium, aur global memory sabse bada par sabse slow (100-1000x zyada latency). Yeh baat samajhna isliye zaroori hai kyunki data ko move karna computation se 100 guna zyada mehenga hota hai.

Ab yahan asli game hai occupancy ka. GPU memory latency ko chhupata hai by switching between threads — matlab jab ek thread memory ka wait kar raha hota hai, GPU dusre thread pe kaam karta hai. Lekin har thread ko kuch registers chahiye hote hain, aur SM ka register budget fixed hota hai (jaise 65,536). Agar har thread zyada registers use karega, toh kam threads ek saath chal paayenge, aur GPU latency ko theek se hide nahi kar paayega — performance gir jaayegi. Isiliye example mein dekha ki 40 registers/thread se 75% occupancy mili, par agar 32 registers/thread kar dein toh 100% occupancy milti hai. Chhoti si register saving bhi bada farak laa sakti hai.

Practical life mein iska matlab yeh hai ki jab aap GPU code likhoge (jaise CUDA), toh aapko soch-samajh ke decide karna padega ki kaunsa data kahaan rakhna hai. Frequently-used data ko registers ya shared memory mein rakho, aur shared memory use karte waqt bank conflicts se bacho taaki access parallel rahe. Yeh intuition aapko fast, efficient GPU programs likhne mein help karega — jo aaj machine learning, graphics, aur scientific computing sab jagah kaam aata hai. Basically, memory hierarchy samajhna hi GPU performance ka asli secret hai.

Memory hierarchy (global, shared, registers)

Overview

Core Concepts

1. Registers

2. Shared Memory

3. Global Memory

Memory Hierarchy Summary Table

Common Mistakes & Misconceptions

Connections

Flashcards

Concept Map

Hinglish (regional understanding)

Hinglish (regional understanding)

Go deeper — visual, from zero

Test yourself — GPU Architecture

Connections