Imagine 32 kids in a class who all have the SAME worksheet but DIFFERENT numbers to plug in. The teacher reads ONE instruction out loud — "now multiply your two numbers" — and all 32 do it at once. That's a warp. If the teacher says "if your name starts with A, do the hard problem, else the easy one," the class must split: A-kids work while everyone watches, then the rest work while A-kids watch — twice as slow! That's divergence. And while one group waits for the librarian (slow memory), the teacher just starts a different group of 32 kids. So nobody sits idle: that's how the GPU stays super busy.
A hardware bundle of 32 threads executed in lockstep from one fetched instruction; the unit of scheduling.
SIMT vs SIMD core difference?
SIMT uses a scalar per-thread programming model with independent per-thread PCs and branching; SIMD uses explicit fixed vector lanes with no per-lane branching.
Formula for global 1D thread index?
i = blockIdx.x * blockDim.x + threadIdx.x
Why add (B-1) when computing number of blocks?
Integer division truncates; (N+B-1)/B rounds up so the last partial block of data isn't dropped.
What is the correct out-of-range guard in a kernel?
if (i >= N) return; — exit only the threads whose index is beyond the data, keeping the valid ones.
What is branch divergence and its cost?
When threads in one warp take different paths, both paths run serialized (masked); cost ≈ T_then + T_else instead of max.
Does divergence break correctness?
No — only performance; each thread still executes its own path correctly via masking.
What does __syncthreads() synchronize?
Only threads within one block (one SM); not across blocks.
What is occupancy?
Active warps on an SM divided by max supported warps; more warps = better latency hiding.
What is memory coalescing?
Combining a warp's 32 addresses into the fewest memory transactions; stride-1 access is ~1 transaction, scattered access is many.
Why choose block size as a multiple of 32?
So warps are fully populated; a non-multiple wastes lanes in the last partial warp.
How does a GPU hide a 400-cycle memory latency?
By switching to another ready warp each cycle (thread oversubscription), keeping ALUs busy.
blockDim vs gridDim?
blockDim = threads per block; gridDim = number of blocks in the grid.
Dekho, GPU ka pura funda ek line mein samajh lo: CPU ek thread ko fast banata hai (bade caches, branch prediction), lekin GPU hazaaron threads chalata hai aur jab koi thread memory ka wait karta hai to turant doosre ready thread pe switch kar deta hai. Isko bolte hain latency hiding — latency avoid nahi karta, usse chhupa deta hai. Isliye GPU mein transistors ka zyada hissa ALUs (calculation units) pe kharch hota hai.
Hardware threads ko 32 ke bundle mein chalata hai jise warp bolte hain. Ek hi instruction fetch hoti hai aur 32 threads usse ek saath, lockstep mein, lekin apne-apne data pe chalate hain. Yahi SIMT hai — tum code aise likhte ho jaise ek scalar thread ke liye (c[i]=a[i]+b[i]), aur hardware khud 32 ko group karke vector ki tarah chalata hai. Index nikalne ka formula simple logic se aata hai: i = blockIdx.x * blockDim.x + threadIdx.x — yaani "apne se pehle ke saare blocks skip karo, phir apne block ke andar ki position add karo." Aur bounds guard hamesha if (i >= N) return; likhna — yaani sirf out-of-range threads ko bahar karo, valid threads ko kaam karne do.
Do sabse bade performance traps yaad rakho. Pehla — warp divergence: agar ek hi warp ke threads if/else mein alag-alag raaste lete hain, to hardware dono raaste serially chalata hai (cost = T_then + T_else). Fix: branch ko warp-uniform banao taki saare 32 ek hi path lein. Doosra — coalescing: agar warp ke threads consecutive memory (a[i]) padhte hain to ek hi transaction mein kaam ho jaata hai (fast); scattered access (a[i*stride]) 32 alag transactions maang sakta hai (slow). Aur ek important baat: divergence sirf speed kharab karta hai, correctness kabhi nahi — har thread apna kaam sahi karta hai, bas masking ki wajah se time lagta hai.
