GPU architecture — SIMT, warps, CUDA model

4.1.25 · Coding › Computer Architecture (Deep)

CUDA thread hierarchy (KYA hai)

Figure — GPU architecture — SIMT, warps, CUDA model

Warps aur SM scheduler (HOW it runs)

Warp divergence (#1 performance trap)

Worked example Divergence worked example

if (threadIdx.x % 2 == 0) heavy_A();   // even lanes
else                       heavy_B();   // odd lanes

Saari 32 lanes ek warp mein hain, aadhi har path leti hain → warp heavy_A run karta hai (16 lanes masked off, kuch nahi karte) PHIR heavy_B. Cost ≈ $T_{A} + T_{B}$ . Fix: poore warps ko agree karaayen, e.g. threadIdx.x / 32 par branch karein ya block-level conditions rakhen taaki har warp uniform ho → cost ≈ $max (T_{A}, T_{B})$ . Yeh step KYUN? Divergence per-warp hoti hai, per-thread nahi; agar warp ke saare 32 threads same path choose karein, toh zero divergence penalty hai.

Memory coalescing (#2 trap)

Common mistakes (Steel-manned)

Recall Feynman: 12-saal ke bachche ko explain karo

Socho 32 bacche ek class mein hain jinke paas ek JAISA worksheet hai lekin ALAG-ALAG numbers plug karne hain. Teacher ek instruction zor se padhta hai — "ab apne dono numbers multiply karo" — aur saare 32 ek saath karte hain. Yeh ek warp hai. Agar teacher bolte hain "agar tumhara naam A se shuru hota hai, mushkil problem karo, warna easy wali," toh class ko split karna padta hai: A-wale kaam karte hain jabki baaki dekhte hain, phir baaki kaam karte hain jabki A-wale dekhte hain — do guna slow! Yeh divergence hai. Aur jab ek group librarian ka intezaar karta hai (slow memory), teacher bas doosre 32 bachchon ka alag group shuru kar deta hai. Toh koi idle nahi baithta: isi tarah GPU super busy rehta hai.

Active-recall flashcards

NVIDIA GPUs par warp kya hota hai?

32 threads ka ek hardware bundle jo ek fetched instruction se lockstep mein execute hota hai; scheduling ki unit.

SIMT vs SIMD ka core difference?

SIMT ek scalar per-thread programming model use karta hai jisme independent per-thread PCs aur branching hoti hai; SIMD explicit fixed vector lanes use karta hai jisme per-lane branching nahi hoti.

Global 1D thread index ka formula?

i = blockIdx.x * blockDim.x + threadIdx.x

Blocks ki sankhya compute karte waqt (B-1) KYUN add karte hain?

Integer division truncate karta hai; (N+B-1)/B round up karta hai taaki data ka last partial block drop na ho.

Kernel mein correct out-of-range guard kya hai?

if (i >= N) return; — sirf un threads ko exit karao jinका index data se bahar hai, valid wale kaam karte rahen.

Branch divergence kya hai aur uski cost?

Jab ek warp ke threads alag paths lete hain, dono paths serialize hokar run karte hain (masked); cost ≈ T_then + T_else instead of max.

Kya divergence correctness tod-ta hai?

Nahi — sirf performance; har thread masking ke zariye apna path sahi execute karta hai.

__syncthreads() kya synchronize karta hai?

Sirf ek block ke andar threads (ek SM); blocks ke across nahi.

Occupancy kya hai?

SM par active warps divided by max supported warps; zyaada warps = better latency hiding.

Memory coalescing kya hai?

Warp ke 32 addresses ko kam se kam memory transactions mein combine karna; stride-1 access ~1 transaction hai, scattered access bahut saare.

Block size 32 ka multiple KYUN choose karein?

Taaki warps fully populated hon; non-multiple last partial warp mein lanes waste karta hai.

GPU 400-cycle memory latency kaise hide karta hai?

Har cycle mein doosre ready warp par switch karke (thread oversubscription), ALUs ko busy rakhta hai.

blockDim vs gridDim?

blockDim = threads per block; gridDim = grid mein blocks ki sankhya.

Connections

SIMD and Vector Instructions (AVX, SSE) — contrast: explicit lanes vs SIMT
CPU Cache Hierarchy and Latency — jo GPUs nahi karte (latency avoid vs hide)
Memory Coalescing and Bandwidth — transaction sizes par deeper dive
CUDA Shared Memory and Tiling — ek block ke andar on-chip memory use karna
Amdahl's and Gustafson's Laws — parallel speedup ki limits
Matrix Multiplication on GPUs — canonical throughput workload