SIMT (single instruction multiple thread)

6.2.4 · Hardware › GPU Architecture

Overview

SIMT (Single Instruction Multiple Thread) NVIDIA ka execution model hai jisme multiple threads same instruction ko lockstep mein execute karte hain, lekin har thread alag data par operate karta hai aur alag execution paths le sakta hai. Yeh SIMD (vector processors) aur true multithreading ke beech ka bridge hai.

Key insight yeh hai: GPUs hazaron similar tasks (pixels, vertices, particles) process karte hain. SIMT throughput maximize karta hai by executing identical instructions across many threads, aur saath mein divergence ko gracefully handle karta hai.

Core Concept: Warps

32 threads per warp kyun?

Hardware efficiency:

• Ek instruction fetch/decode unit 32 execution units ko serve karta hai
• 32 lanes of ALUs parallel mein execute karte hain
• Memory coalescing 32-thread aligned access patterns ke saath best kaam karta hai
• 32 power of 2 hai, jo address calculations simplify karta hai (thread_id & 31 lane within warp deta hai)

Warp utilization ka derivation: $\text{Warp Utilization}=\frac{\text{Active Threads}}{\text{Warp Size}}=\frac{N_{active}}{32}$

Agar tum 100 threads ke saath kernel launch karo:

• 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 waste!

Yeh kyun matter karta hai: Execution lanes waste na ho isliye hamesha thread counts launch karo jo 32 ke multiples hon.

SIMT vs SIMD: Critical Difference

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

• Divergence impossible—saare lanes MUST execute karte hain
• Software vectorization explicitly manage karta hai
• Example: CPUs par SSE/AVX

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

• Har thread ke paas apne registers aur program counter hote hain (logical view)
• Divergence allowed via predication masks
• Programmer scalar code likhta hai; hardware threads ko broadcast karta hai

Derivation: SIMT ko predication kyun chahiye

Jab warp ke threads alag branches lete hain:

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

Kya hota hai: Warp dono branches serially execute karta hai:

1. Phase 1: Threads 0,2,4,... active, computeA() execute karte hain. Threads 1,3,5,... masked off.
2. Phase 2: Threads 1,3,5,... active, computeB() execute karte hain. 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: Saare 32 threads unique paths lete hain → 32× slowdown! Isliye branch divergence GPU performance ko barbad kar deta hai.

Scenario: 16 threads hot path par, 16 cold path par jaate hain.

Divergence ke bina time: 10 cycles (agar sab hot) ya 30 cycles (agar sab cold) 50/50 split ke saath time: 10 + 30 = 40 cycles Efficiency: 16/(32) = 50% warp utilization hot path phase ke dauran, 50% cold path phase ke dauran

Hardware Implementation

Har SM mein hota hai:

• Warp Scheduler: Ready warp select karta hai, instruction issue karta hai
• 32 CUDA Cores (FP32 ALUs): Threads across instruction execute karte hain
• Predication Masks: 32-bit register, 1 bit per thread (1=active, 0=masked)

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

Modern GPUs (Ampere): SM per 4 warp schedulers → Per cycle 4 instructions issue kar sakte hain (4 alag warps ko)

Warp scheduling se latency hide karna

Problem: Memory access ~400 cycles leta hai. GPU stall kyun nahi hota?

SIMT solution: Jab Warp 0 memory ka wait kar raha hota hai, Warp 1, Warp 2, etc. par switch kar lo.

Occupancy requirement ka derivation:

Maano:

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

Latency $L$ hide karne ke liye, tumhe chahiye: $W\ge L\cdot T$

Example: 400-cycle memory latency hide karna, 1 instruction/cycle chahiye → $W\ge 400\cdot 1=400$ warps worth of instructions chahiye

Agar har warp average par har 4 cycles mein issue karta hai: → $W\ge 400/4=100$ warps per SM

Lekin SMs typically sirf 32-64 warps max support karte hain! Isliye:

1. GPUs mein bahut saare SMs hote hain (sirf deep pipelines nahi)
2. Kernels mein high occupancy honi chahiye (bahut saare active warps) latency hide karne ke liye

SM mein 4 active warps hain, har ek alag stage par:

• Warp 0: Memory ka wait kar raha hai (cycle 1/400)
• Warp 1: Ready, ADD execute kar raha hai
• Warp 2: Memory ka wait kar raha hai (cycle 300/400)
• Warp 3: Ready, MUL execute kar raha hai

Cycle N: Scheduler Warp 1 pick karta hai, 32 cores ko ADD issue karta hai Cycle N+1: Scheduler Warp 3 pick karta hai, 32 cores ko MUL issue karta hai Cycle N+2: Warp 0 abhi ready nahi, Warp 2 abhi ready nahi, scheduler Warp 1 replay kar sakta hai ya stall ho sakta hai

Key insight: SIMT latency massive multithreading se tolerate karta hai, speculation ya out-of-order execution se nahi (jo CPUs use karte hain).

Thread Hierarchy aur SIMT

Hardware se mapping: $\text{Block size}\to \text{\# Warps}=\lceil \frac{\text{Block size}}{32}\rceil$

Kernel launch karo:

• Block size: 256 threads
• 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

Agar GPU mein 80 SMs hain, har ek 16 warps concurrently run karta hai: Concurrent warps: $80\times 16=1280$ warps Waves: $\lceil 8000/1280\rceil =7$ waves kernel complete karne ke liye

Har wave mein time $T_{wave}$ lagta hai, total time $\approx 7\cdot T_{wave}$ (overlap ignore karke).

SIMT ke liye Optimize karna

1. Divergence se bachao

Divergence cost quantify karna: Agar warp ka $f$ fraction $N$ branches mein diverge karta hai, execution time: $T={\sum }_{i=1}^N{T}_i$

vs. converged: $T={\max }_i(T_i)$

Best case: Saare threads same path → $N=1$ Worst case: Har thread unique path → $N=32$

Fix strategies:

• Data reorganize karo taaki similar threads same warp mein grouped hon
• Branches avoid karne ke liye warp-level primitives (__ballot_sync, __shfl_sync) use karo
• Processing se pehle data ko branch outcome ke hisaab se sort karo

2. Memory coalescing

Har warp ek memory transaction issue karta hai. Agar 32 threads addresses access karte hain: ${\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% (128 bytes transfer, 128 bytes use)
• Strided 32: 3.125% (4096 bytes transfer, 128 bytes use)
• Random: ~3% (likely 32 separate cache lines)

3. Occupancy tuning

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

High occupancy kyun help karta hai: $\text{Latency hiding capacity}\propto \text{\# Active Warps}$

Occupancy limited hota hai:

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

GPU: 65536 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% ✓

Agar kernel 128 registers/thread use kare: Threads per SM = $\lfloor 65536/128\rfloor =512$ threads → Occupancy = 512/1024 = 50% (wasted potential!)

Common Mistakes

Yeh sahi kyun lagta hai: Dono parallel mein multiple data elements par same instruction execute karte hain.

Yeh galat kyun hai:

1. Thread independence: SIMT threads ke paas alag registers hote hain aur woh diverge kar sakte hain. SIMD vector lanes nahi kar sakte—divergence ke liye software ko masks se emulate karna padta hai.
2. Programming model: SIMT tumhe scalar code likhne deta hai (ek thread ki logic). SIMD mein explicit vectorization chahiye.
3. Hardware scheduling: SIMT scheduler dynamically hazaron threads manage karta hai. SIMD sirf ek single thread ke instruction stream se vector instructions execute karta hai.

Fix: SIMT ek hardware implementation hai jo parallel execution ko programmer ko multithreading jaisa dikhata hai, automatic broadcasting aur divergence handling ke saath.

Yeh sahi kyun lagta hai: Mujhe sirf 100 computations chahiye, toh 100 threads efficient hone chahiye.

Yeh galat kyun hai:

• 100 threads = 4 warps (32+32+32+4)
• Last warp: 4/32 = 12.5% utilized, 28 ALUs idle
• Agar yahi ek block hai, toh tum poore GPU par sirf 4 warps use kar rahe ho (80 SMs × 64 warps/SM = 5120 warp slots available!) → 0.08% GPU utilization

Fix: 32 ke multiples tak round up karo. Chhote kaam ke liye, block size badhao ya GPU saturate karne ke liye zyada blocks launch karo. 100 threads ko 128 threads (4 full warps) hona chahiye many blocks ki grid mein.

Yeh sahi kyun lagta hai: Zyada active warps = better latency hiding = faster execution.

Yeh galat kyun hai:

• Memory bandwidth: Agar kernel memory-bound hai, toh zyada warps add karne se same bandwidth ke liye sirf zyada congestion hogi.
• Shared memory: Zyada blocks fit karne ke liye shared memory reduce karna global memory accesses force kar sakta hai, jisse higher occupancy ke bawajood cheezein slow ho jaati hain.
• Instruction mix: Compute-heavy kernels occupancy se benefit karte hain. Memory-heavy kernels aksar nahi karte.

Fix: Occupancy tab tak optimize karo jab tak "kaafi achha" na ho (>50%), phir actual bottleneck par focus karo (bandwidth, divergence, instruction throughput). Profiler use karo dekhne ke liye ki kya occupancy badhane se performance improve hoti hai.

Yeh sahi kyun lagta hai: SIMT divergence ko predication ke saath gracefully handle karta hai—threads alag paths le sakte hain!

Yeh galat kyun hai: "Gracefully handle karna" ≠ "free". Divergence execution serialize karta hai: $T_{diverged}=\sum \text{(time for each unique path)}$

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

Fix: SIMT divergence allow karta hai lekin penalize bhi karta hai. Algorithms design karo divergence minimize karne ke liye: data sort karo, predication use karo (result = condition ? a : b branchless select instruction mein compile ho sakta hai), ya dono branches compute karo aur result select karo.

Feynman Technique

Lekin yahan twist hai: Kya hoga agar sawaal 5 poochhe "Agar tumhara birthday saal ke pehle aadhe mein hai, toh 2 se multiply karo. Warna, 10 add karo." Ab aadhe students multiply karte hain, aadhe add karte hain. Tum dono instructions ek saath nahi padh sakte!

Toh tum kehte ho: "January-June birthday wale, abhi apni calculation karo. Baaki sabko wait karo." Phir: "July-December birthday wale, ab tumhari baari. Baaki wait karo."

Problem: Ab isme DOUBLE time lagta hai kyunki tumhe class ko do groups mein split karna pada jo alag cheezein kar rahe the. Ise branch divergence kehte hain.

GPU ek teacher ki tarah hai jo ek saath hazaron classrooms (warps) manage karta hai. Har classroom mein 32 students (threads) hain jo saath mein follow karte hain. Agar students alag-alag cheezein karte rahe, poora system slow ho jaata hai. Lekin agar sab sync mein hain, toh tum hazaron students ko 32 students padhane se thodi zyada mehnat mein padha sakte ho!

Lesson: Similar kaam ko saath group karo (same calculation, same if/else path) taaki tumhare GPU classrooms sync mein rahen.

Socho: "SIMT = SIMultaneous, lekin Threads diverge ho sakte hain" vs "SIMD = SIMultaneous, Data locked together"

Ya: "32 workers, 1 instruction broadcast, Threads breaks le sakte hain"

Connections

• SIMD vs SIMT Comparison - contrasting execution models
• GPU Memory Coalescing - SIMT mein aligned memory access kyun chahiye
• Warp Divergence - thread divergence ki performance cost
• CUDA Thread Hierarchy - threads/warps/blocks hardware se kaise map hote hain
• GPU Occupancy - latency hiding ke liye concurrent warps maximize karna
• Streaming Multiprocessor - warps execute karne wali hardware unit
• Instruction-Level Parallelism - CPU approach vs GPU approach
• GPU vs CPU Architecture - SIMT throughput-oriented design mein kyun fit hota hai

#flashcards/hardware

SIMT kya hai? :: Single Instruction Multiple Thread—NVIDIA ka execution model jisme warp ke 32 threads lockstep mein same instruction execute karte hain lekin alag data par operate kar sakte hain aur alag execution paths par diverge ho sakte hain.

32 threads per warp kyun? :: Hardware efficiency—ek instruction fetch/decode 32 execution units serve karta hai; memory coalescing 32-thread aligned access ke saath best kaam karta hai (power of 2); parallel execution aur divergence overhead ke beech balance karta hai.

Concept Map