SIMT (single instruction multiple thread)

Overview

SIMT (Single Instruction Multiple Thread) is NVIDIA's execution model where multiple threads execute the same instruction in lockstep, but each thread operates on different data and can take different execution paths. It bridges SIMD (vector processors) and true multithreading.

The key insight: GPUs process thousands of similar tasks (pixels, vertices, particles). SIMT maximizes throughput by executing identical instructions across many threads while handling divergence gracefully.

Core Concept: Warps

Why 32 threads per warp?

Hardware efficiency:

• One instruction fetch/decode unit serves32 execution units
• 32 lanes of ALUs execute in parallel
• Memory coalescing works best with 32-thread aligned access patterns
• 32 is a power of 2, simplifying address calculations (thread_id & 31 gives lane within warp)

Derivation of warp utilization: $\text{Warp Utilization} = \frac{\text{Active Threads}}{\text{Warp Size}} = \frac{N_{active}}{32}$

If you launch a kernel with 100 threads:

• Warp 0: 32 threads (100% utilized)
• Warp 1: 32 threads (100% utilized)
• Warp 2: 32 threads (100% utilized)
• Warp 3: 4 threads (12.5% utilized) ← 28 lanes wasted!

Why this matters: Always launch thread counts that are multiples of 32 to avoid wasting execution lanes.

SIMT vs SIMD: The Critical Difference

SIMD (traditional vector processors): $\text{One instruction} \rightarrow \text{N data elements in vector register}$

• Divergence impossible—all lanes MUST execute
• Software manages vectorization explicitly
• Example: SSE/AVX on CPUs

SIMT (GPU model): $\text{One instruction} \rightarrow \text{N independent threads}$

• Each thread has its own registers and program counter (logical view)
• Divergence allowed via predication masks
• Programmer writes scalar code; hardware broadcasts to threads

Derivation: Why SIMT needs predication

When threads in a warp take different branches:

if (threadIdx.x % 2 == 0) {
    result = computeA();  // Branch A
} else {
    result = computeB();  // Branch B
}

What happens: The warp executes both branches serially:

1. Phase 1: Threads 0,2,4,... active, execute computeA(). Threads 1,3,5,... masked off.
2. Phase 2: Threads 1,3,5,... active, execute computeB(). Threads 0,2,4,... masked off.

Time cost: $T_{diverged} = T_A + T_B$ vs. fully converged: $T_{converged} = \max(T_A, T_B)$

Worst case: All32 threads take unique paths → 32× slowdown! This is why branch divergence kills GPU performance.

Scenario: 16 threads go to hot path, 16 to cold path.

Time without divergence: 10 cycles (if all hot) or 30 cycles (if all cold)
Time with 50/50 split: 10 + 30 = 40 cycles
Efficiency: 16/(32) = 50% warp utilization during hot path phase, 50% during cold path phase

Hardware Implementation

Each SM has:

• Warp Scheduler: Selects ready warp, issues instruction
• 32 CUDA Cores (FP32 ALUs): Execute instruction across threads
• Predication Masks: 32-bit register, 1bit per thread (1=active, 0=masked)

Instruction dispatch rate: $\text{Instructions per cycle} = \frac{\text{\# Warp Schedulers}}{\text{Cycles per instruction}}$

Modern GPUs (Ampere): 4 warp schedulers per SM
→ Can issue 4 instructions per cycle (to 4 different warps)

Hiding latency via warp scheduling

The problem: Memory access takes ~400 cycles. Why doesn't the GPU stall?

SIMT solution: While Warp 0 waits for memory, switch to Warp 1, Warp 2, etc.

Derivation of occupancy requirement:

Let:

• $L$ = latency of operation (cycles)
• $W$ = warps per SM
• $T$ = throughput requirement (instructions/cycle)

To hide latency $L$, you need: $W \geq L \cdot T$

Example: Hide 400-cycle memory latency, need 1 instruction/cycle
→ Need $W \geq 400 \cdot 1 = 400$ warps worth of instructions

If each warp issues every 4 cycles on average:
→ $W \geq 400/4 = 100$ warps per SM

ButMs typically support 32-64 warps max! This is why:

1. GPUs have many SMs (not just deep pipelines)
2. Kernels must have high occupancy (many active warps) to hide latency

SM has 4 active warps, each at different stages:

• Warp 0: Waiting for memory (cycle 1/400)
• Warp 1: Ready, executing ADD
• Warp 2: Waiting for memory (cycle 300/400)
• Warp 3: Ready, executing MUL

Cycle N: Scheduler picks Warp 1, issues ADD to32 cores
Cycle N+1: Scheduler picks Warp 3, issues MUL to 32 cores
Cycle N+2: Warp 0 still not ready, Warp 2 still not ready, scheduler might replay Warp 1 or stall

Key insight: SIMT tolerates latency through massive multithreading, not speculation or out-of-order execution (which CPUs use).

Thread Hierarchy and SIMT

Mapping to hardware: $\text{Block size} \rightarrow \text{\# Warps} = \lceil \frac{\text{Block size}}{32} \rceil$

Launch kernel with:

• Block size: 256threads
• Grid size: 1000 blocks

Per block: $256 / 32 = 8$ warps
Total warps: $1000 \times 8 = 8000$ warps
Total threads: $1000 \times 256 = 256{,}000$ threads

If GPU has 80 SMs, each running 16 warps concurrently:
Concurrent warps: $80 \times 16 = 1280$ warps
Waves: $\lceil 8000/1280 \rceil = 7$ waves to complete kernel

Each wave takes time $T_{wave}$, total time $\approx 7 \cdot T_{wave}$ (ignoring overlap).

Optimizing for SIMT

1. Avoid divergence

Quantifying divergence cost: If $f$ fraction of warp diverges into $N$ branches, execution time: $T = \sum_{i=1}^{N} T_i$

vs. converged: $T = \max_i(T_i)$

Best case: All threads same path → $N=1$
Worst case: Every thread unique path → $N=32$

Fix strategies:

• Reorganize data so similar threads are grouped in same warp
• Use warp-level primitives (__ballot_sync, __shfl_sync) to avoid branches
• Sort data by branch outcome before processing

2. Memory coalescing

Each warp issues one memory transaction. If 32 threads access addresses: $\text{addr}_i = \text{base} + i \cdot \text{stride}$

Coalesced (stride=4 bytes for int): 1 transaction (128 bytes)
Strided (stride=128 bytes): Up to 32 transactions!

Efficiency: $\text{Coalescing Efficiency} = \frac{\text{Bytes Requested}}{\text{Bytes Transferred}}$

// GOOD: Coalesced (stride = 4)
int val = data[threadIdx.x];  
 
// BAD: Strided (stride = 32*4 = 128)
int val = data[threadIdx.x * 32];
 
// TERRIBLE: Random
int val = data[random_indices[threadIdx.x]];

Bandwidth utilization:

• Coalesced: 100% (transfer 128 bytes, use 128 bytes)
• Strided 32: 3.125% (transfer 4096 bytes, use 128 bytes)
• Random: ~3% (likely 32 separate cache lines)

3. Occupancy tuning

Occupancy = (Active warps) / (Max warps per SM)

Why high occupancy helps: $\text{Latency hiding capacity} \propto \text{\# Active Warps}$

Occupancy limited by:

• Registers per thread: $\text{Warps} \leq \frac{\text{Total Registers}}{\text{Registers per thread} \times 32}$
• Shared memory per block: $\text{Blocks per SM} \leq \frac{\text{Total Shared Mem}}{\text{Shared Mem per Block}}$
• Max threads per SM: $\text{Warps} \leq \frac{\text{Max Threads per SM}}{32}$

GPU: 6536 registers/SM, 1024 threads/SM max
Kernel: 64 registers/thread, block size 256

Register limit:
Threads per SM = $\lfloor 65536 / 64 \rfloor = 1024$ threads
Warps = $1024 / 32 = 32$ warps

Thread limit:
Max warps = $1024 / 32 = 32$ warps

Block limit:
Blocks per SM = $\lfloor 1024 / 256 \rfloor = 4$ blocks
Warps = $4 \times (256/32) = 32$ warps

Achieved occupancy:32/32 = 100% ✓

If kernel used 128 registers/thread:
Threads per SM = $\lfloor 6536 / 128 \rfloor = 512$ threads
→ Occupancy = 512/1024 = 50% (wasted potential!)

Common Mistakes

Why this feels right: Both execute same instruction across multiple data elements in parallel.

Why it's wrong:

1. Thread independence: SIMT threads have separate registers and can diverge. SIMD vector lanes cannot—divergence requires software to emulate with masks.
2. Programming model: SIMT lets you write scalar code (one thread's logic). SIMD requires explicit vectorization.
3. Hardware scheduling: SIMT scheduler manages thousands of threads dynamically. SIMD just executes vector instructions from a single thread's instruction stream.

The fix: SIMT is a hardware implementation that makes parallel execution look like multithreading to the programmer, with automatic broadcasting and divergence handling.

Why this feels right: I only need 100 computations, so 100 threads should be efficient.

Why it's wrong:

• 100 threads = 4 warps (32+32+32+4)
• Last warp: 4/32 = 12.5% utilized, 28 ALUs idle
• If this is the only block, you're using 4 warps across the entire GPU (80 SMs × 64 warps/SM = 5120 warp slots available!) → 0.08% GPU utilization

The fix: Round up to multiples of 32. For small work, increase block size or launch more blocks to saturate the GPU.100 threads should be 128 threads (4 full warps) in a grid of many blocks.

Why this feels right: More active warps = better latency hiding = faster execution.

Why it's wrong:

• Memory bandwidth: If kernel is memory-bound, adding more warps just creates more congestion fighting for the same bandwidth.
• Shared memory: Reducing shared memory to fit more blocks might force more global memory accesses, slowing things down despite higher occupancy.
• Instruction mix: Compute-heavy kernels benefit from occupancy. Memory-heavy kernels often don't.

The fix: Optimize occupancy until it's "good enough" (>50%), then focus on the actual bottleneck (bandwidth, divergence, instruction throughput). Use profiler to see if increasing occupancy improves performance.

Why this feels right: SIMT handles divergence gracefully with predication—threads can take different paths!

Why it's wrong: "Handles gracefully" ≠ "free". Divergence serializes execution: $T_{diverged} = \sum \text{(time for each unique path)}$

With2 branches equally likely: $T = 2 \times T_{branch}$ (2× slowdown).

The fix: SIMT allows divergence but penalizes it. Design algorithms to minimize divergence: sort data, use predication (result = condition ? a : b can compile to a branchless select instruction), or compute both branches and select result.

Feynman Technique

But here's the twist: What if question 5 asks "If your birthday is in the first half of the year, multiply by 2. Otherwise, add 10." Now half the students multiply, half add. You can't read both instructions at once!

So you say: "Those with birthdays January-June, do your calculation now. Everyone else, wait." Then: "July-December birthdays, now it's your turn. Others, wait."

The problem: Now it takes TWICE as long because you had to split the class into two groups doing different things. That's called branch divergence.

The GPU is like a teacher managing thousands of classrooms (warps) at once. Each classroom has32 students (threads) who must follow along together. If students keep doing different things, the whole system slows down. But if everyone's in sync, you can teach thousands of students with barely more effort than teaching32!

The lesson: Group similar work together (same calculation, same if/else path) so your GPU classrooms stay in sync.

Think: "SIMT = SIMultaneous, but Threads can diverge"
vs "SIMD = SIMultaneous, Data locked together"

Or: "32 workers, 1 instruction broadcast, Threads can take breaks"

Connections

• SIMD vs SIMT Comparison - contrasting execution models
• GPU Memory Coalescing - why SIMT requires aligned memory access
• Warp Divergence - the performance cost of thread divergence
• CUDA Thread Hierarchy - how threads/warps/blocks map to hardware
• GPU Occupancy - maximizing concurrent warps for latency hiding
• Streaming Multiprocessor - the hardware unit executing warps
• Instruction-Level Parallelism - CPU approach vs GPU approach
• GPU vs CPU Architecture - why SIMT fits throughput-oriented design

#flashcards/hardware

What is SIMT? :: Single Instruction Multiple Thread—NVIDIA's execution model where32 threads in a warp execute the same instruction in lockstep but can operate on different data and diverge to different execution paths.

Why32 threads per warp? :: Hardware efficiency—one instruction fetch/decode serves 32 execution units; memory coalescing works best with 32-thread aligned access (power of 2); balances parallel execution with divergence overhead.

Concept Map

Hinglish (regional understanding)