Streaming multiprocessors (SM)

Overview

The Streaming Multiprocessor (SM) is the fundamental execution unit inside a GPU. Think of it as a mini-processor optimized for parallel execution – while a CPU might have 4-16 cores, a modern GPU has dozens of SMs, and each SM can execute hundreds of threads simultaneously.

This is hardware-level parallelism: multiple SMs run completely independently, each jugling hundreds of threads. The GPU scheduler distributes thread blocks across SMs – one SM might render pixels0-1023 while another renders 1024-2047.

Architecture of a Single SM

What's Inside an SM?

Why this structure? Each component targets a specific bottleneck:

• CUDA cores → raw arithmetic throughput
• SFUs → offload expensive math (20-100 cycles on CPU, 1-2 cycles on SFU)
• Registers → eliminate memory latency (threads keep variables in registers)
• Shared memory → fast inter-thread communication (100x faster than global memory)

Hierarchical Execution Model

The SM executes threads in a strict hierarchy:

Why 32 threads per warp? This is SIMT (Single Instruction Multiple Thread) – the warp scheduler issues one instruction that all 32 threads execute simultaneously. It's like SIMD (vector processing) but with more flexibility: threads can take different branches, though it costs performance (divergence).

Derivation of thread scheduling:

Each cycle, the warp scheduler:

1. Picks an eligible warp (threads not stalled on memory/dependencies)
2. Issues an instruction to that warp's 32 threads
3. The CUDA cores execute the instruction in parallel

If a warp stalls (waiting for memory), the scheduler immediately switches to another warp – this is why GPUs need thousands of threads: to hide latency.

Occupancy calculation:

Why occupancy matters: Higher occupancy → more warps available → better latency hiding. But occupancy is limited by:

• Registers per thread × threads per block ≤ Total registers
• Shared memory per block ≤ Total shared memory
• Blocks per SM ≤ Max blocks (usually 16-32)

Example: If your kernel uses 64 registers/thread and the SM has 65,536 registers, max threads = 65,536 / 64 = 1024. But if max warps = 64, you cap at 64 × 32 = 2048 threads anyway.

How Work Gets Distributed

Scenario: Launch a kernel with 10,000 threads organized as 40 blocks of 256 threads each. GPU has 80 SMs.

What happens:

1. Block assignment: The GPU's Grid Distribution Unit assigns blocks to SMs. Each block runs on exactly one SM from start to finish (no migration).

   • With 80 SMs and 40 blocks, each SM gets ⌊40/80⌋ = 0 or 1 block initially, then receives more as blocks complete.

2. Warp creation: Each 256-thread block becomes 256/32 = 8 warps. The SM's warp scheduler now juggles these 8 warps.

3. Instruction issue: Each cycle, the scheduler picks one eligible warp and issues its next instruction to the CUDA cores.

   • If Warp 0 is fetching from memory (300 cycles), the scheduler runs other eligible warps while Warp 0's data is in flight.

Why this step? (Latency-hiding math done correctly) Suppose a global-memory access stalls a warp for 300 cycles, and a scheduler issues 1 instruction per cycle. To fully hide that stall you need roughly 300 independent in-flight warps (one per cycle of latency) — one SM's 8 warps from a single block are not enough to cover a 300-cycle stall alone. In practice we approach this by having many blocks resident at once: with a real SM cap of ~48-64 warps, plus multiple schedulers each issuing per cycle, we get closer to hiding hundreds of cycles. The lesson: maximize resident warps, because 8 warps only cover ~8 cycles of latency per scheduler pass.

Problem: Your kernel uses 128 registers per thread. SM has 65,536 registers. Block size = 256 threads.

Calculation:

• Registers needed per block: 256 × 128 = 32,768
• Max blocks per SM: 65,536 / 32,768 = 2 blocks
• Warps per block: 256/32 = 8 → Total warps = 2 × 8 = 16 warps
• If max warps/SM = 64, occupancy = 16/64 = 25%

Why this step? High register usage starves the SM of parallelism. With only 16 warps, and one scheduler issuing ~1 instruction/warp/cycle, you cover only a small slice of a 400-800 cycle global-memory stall — so a memory-bound kernel will stall. To hide hundreds of cycles you want tens of resident warps, which register pressure here prevents.

Fix: Reduce registers (algorithmic change), increase resident blocks/warps, or accept lower occupancy if compute-bound.

Memory Hierarchy and SM

Each SM has multiple memory tiers, each with distinct latency characteristics:

Note: Shared memory and the L1 data cache often share the same physical SRAM but are not the same thing — shared memory is explicitly managed (you index into it directly, no tag lookup), while L1 is hardware-managed with tag checks and eviction. Their latencies and behavior therefore differ by architecture.

Why shared memory? Threads in a block can cooperate via shared memory:

__shared__ float cache[256];  // Visible to all threads in block
cache[threadIdx.x] = data[i]; // Each thread writes
__syncthreads();              // Wait for all
float sum = cache[0] + cache[1] + ..; // Read others' data

Shared memory is ~100x faster than global memory and enables algorithms like tiled matrix multiply, reductions, and prefix sums.

Bank conflicts (shared memory only): Shared memory is divided into 32 banks. If multiple threads in a warp access different addresses in the same bank, the accesses serialize. Design access patterns so thread i accesses bank i % 32. (Registers are private per-thread and do not suffer bank conflicts in this way.)

Warp Divergence

Why it feels right: The GPU marketing says "thousands of threads in parallel!"

Reality: Threads in a warp execute the same instruction. If threads take different branches:

if (threadIdx.x < 16) {
    A;  // Threads 0-15
} else {
    B;  // Threads 16-31
}

The warp must execute both paths serially:

1. Execute A with threads 0-15 active (16-31 masked)
2. Execute B with threads 16-31 active (0-15 masked)

Performance cost: 2× the instructions. For complex branches, even worse.

Steel-man: The GPU does allow divergence for flexibility (unlike pure SIMD). But it's not free – use it sparingly.

Fix: Reorganize data so threads in a warp take the same branch. Example: sort data by category before processing.

You divide the 10,000 squares into groups of 256 (blocks) and give each helper a group. The helper further splits their 256 squares into 8 mini-groups of 32 (warps). Every second, they pick one mini-group and color all 32 squares with the same crayon stroke.

If a mini-group is waiting for a new crayon box to arrive (memory fetch), the helper switches to another mini-group immediately – no wasted time! But a crayon box can take a long time to arrive (like 300 seconds), so you actually need lots of mini-groups waiting so the helper always has something to color. That's why you need thousands of threads: to fully cover the long waits.

And here's a trick: if your instructions say "color squares 0-15 red and 16-31 blue," the helper has to do two strokes (red, then blue) instead of one. That's slower! So it's best when all 32 squares in a mini-group need the same color.

Or: Streaming Machine – streams of threads flow through it continuously.

Connections

• CUDA-cores – The arithmetic units inside each SM
• Warp-scheduling – How the SM picks which threads run each cycle
• Shared-memory – Fast inter-thread communication within an SM
• Thread-blocks – The unit of work assigned to an SM
• Occupancy-optimization – Tuning resource usage to maximize active warps
• GPU-memory-hierarchy – How SMs access L1/L2/global memory
• Tensor-cores – Specialized matrix units in modern SMs

#flashcards/hardware

Concept Map

Hinglish (regional understanding)

Hinglish (regional understanding)