CUDA cores and execution model — Hinglish notes

6.2.3 · Hardware › GPU Architecture

CUDA core kya hota hai?

Key distinction: Ek CPU core ka apna instruction decoder, branch predictor, L1/L2 cache hota hai, aur ye independent instruction streams chala sakta hai. Ek CUDA core apni warp ke 31 doosre cores ke saath control logic share karta hai (ek group of 32 threads jo lockstep mein execute hote hain).

Architecture hierarchy (neeche se upar)

CUDA core (scalar ALU): 1 FP32/integer op per cycle
Warp (execution unit): 32 threads ek hi instruction execute karte hain
Streaming Multiprocessor (SM): 64-128 CUDA cores contain karta hai (architecture dependent), shared memory, warp schedulers, registers
GPU chip: 10-100+ SMs (e.g., RTX 4090 mein 128 SMs × 128 cores/SM = 16,384 CUDA cores hain)

SIMT execution model

Derivation: Warp mein 32 threads kyun?

Design constraint: SM ek cycle mein ek warp ko ek instruction issue karta hai. Warp scheduler ko ye karna hota hai:

Memory latency hide karna (global memory ke liye 100-400 cycles)
ALUs ko kaam se fed rakhna
Register file constraints mein fit karna

Warp cores par kaise map hota hai: Jab tak SM mein kam se kam utne ALUs hain jitne ek warp mein threads hain (yaani ≥ 32), ek single warp instruction ek cycle mein complete hoti hai; koi extra cores us instruction ke liye simply idle rehte hain (unhe doosre warps use karte hain). Toh:

Ek SM with 32 cores → ek warp instruction per cycle, saari cores busy.
Ek SM with 64 ya 128 cores → ek warp instruction phir bhi ek cycle mein complete hoti hai, lekin SM ek cycle mein multiple warp instructions issue kar sakta hai (alag warps se, multiple schedulers ke zariye) extra cores ko busy rakhne ke liye.

32 ka warp size ek hardware convention hai (2 ki power jo address masking aur indexing simplify karta hai, aur 8/16/32 ki SIMD-width lineage se match karta hai). Ye kisi single SM ke core count se derived nahi hai.

Register pressure (warp size occupancy ke saath kyun interact karta hai): SM per register file architecture-dependent hoti hai (commonly 65,536 32-bit registers kaafi recent NVIDIA SMs par, lekin ye vary karta hai). Occupancy total registers se limited hoti hai: agar SM $M$ registers par cap kare aur aap $P$ resident threads chahein, toh har thread approximately $\frac{M}{P}$ registers use kar sakta hai. For example, $M = 65, 536$ aur $P = 2048$ resident threads ke saath, ye roughly $\frac{65536}{2048} = 32$ registers/thread hai. Ye ek rough ceiling hai — real allocation per warp aur per architecture quantized hoti hai, toh 32 ko ek illustrative figure mano, universal law nahi.

Thread hierarchy aur scheduling

Blocks kyun? Ye ek scalability abstraction provide karte hain. 1,000,000 threads ka ek kernel 1000 blocks of 1000 threads mein organize hoke 10 SMs (100 blocks/SM) ya 80 SMs (~12 blocks/SM) wale GPU par bina code changes ke chal sakta hai.

Warp scheduling mechanism

Har SM mein 1-4 warp schedulers hote hain. Har cycle, ek scheduler:

Ek ready warp select karta hai (memory/dependency par stalled nahi)
Uski next instruction CUDA cores ko issue karta hai
Next cycle mein next warp par move karta hai (zero-overhead context switch, kyunki saare warps ke registers pehle se loaded hain)

Latency hiding derivation (carefully kiya hua):

Bookkeeping setup karte hain. Maan lo:

$S$ = SM par warp schedulers ki sankhya (har ek ek warp-instruction per cycle issue karta hai),
ek warp $C$ cycles ka useful compute karta hai aur phir ek long-latency memory op issue karta hai jo use $L$ cycles ke liye stall karta hai,
$T$ = resident (schedulable) warps ki sankhya.

Jab ek warp $L$ cycles ke liye stalled hai, hum chahte hain schedulers ke paas issue karne ke liye doosre ready warps hon. Har cycle SM $S$ warp-instructions issue kar sakta hai, toh $L$ stall cycles mein SM mein $S \cdot L$ warp-instructions issue karne ki capacity hai. Us window mein ALUs ko busy rakhne ke liye, hume itne independent warps chahiye jo itni instructions supply karein. Kyunki har warp stall hone se pehle approximately $C$ compute instructions supply karta hai, latency fully cover karne ke liye zaruri warps ki sankhya hai:

$T_{required} = \frac{S \cdot L}{C} .$

Worked numbers: $S = 4$ , $L = 200$ , $C = 10$ ke saath: $T_{required} = \frac{4 \times 200}{10} = 80 warps .$

Lekin ek typical SM ~64 resident warps tak cap karta hai (2048 threads ÷ 32). Kyunki $80 > 64$ , full latency hiding yahan impossible hai — exactly isliye memory-efficient kernels (coalescing, caching, higher compute intensity $C$ ) matter karte hain: $C$ badhane se $T_{required}$ kam hota hai.

Memory hierarchy aur CUDA cores

CUDA cores data ek hierarchy ke through access karte hain:

Registers (per-thread, ~1 cycle latency): Sabse fast, ~32-64 per thread tak limited
Shared memory (per-block, ~30 cycle latency): Explicitly managed, 48-96 KB/SM
L1 cache (per-SM, ~30 cycles): Implicit, recent architectures mein shared mem ke saath combined
L2 cache (global, ~200 cycles): Saare SMs mein shared, 6-40 MB
Global memory (DRAM, ~400 cycles): 8-24 GB, high bandwidth (1 TB/s) lekin high latency

Coalesced memory access: Jab ek warp ke threads consecutive memory addresses access karte hain (e.g., thread 0 → A[0], thread 1 → A[1], ..., thread 31 → A[31]), hardware unhe ek single 128-byte transaction mein combine kar deta hai. Uncoalesced access (random addresses) → up to 32 separate transactions → up to 32× slower.

Worked example Simple vector addition kernel

__global__ void vectorAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}
 
// Launch: <<<numBlocks, threadsPerBlock>>>
int threadsPerBlock = 256;
int numBlocks = (N + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, N);

Hood ke neeche kya hota hai:

Grid launch: GPU blocks ko SMs mein distribute karta hai (e.g., N=1,000,000 → 3907 blocks @ 256 threads)
Block → SM mapping: Har SM ko multiple blocks milte hain (registers, shared mem se limited)
Warp creation: Har block ke 256 threads → 8 warps (256 ÷ 32)
Execution: Warp 0 ke saare 32 threads compute karte hain i = blockIdx.x * 256 + [0...31], A[i] aur B[i] parallel mein load karte hain (agar memory access coalesced hai), add karte hain, C[i] store karte hain
Warp 1 next execute hota hai threadIdx.x = [32...63] ke saath, aur aise aage Ye step kyun? if (i < N) isliye zaroori hai kyunki N necessarily 256 se divisible nahi ho sakta. Ek warp ke andar, kuch threads (i ≥ N) inactive honge, lekin warp phir bhi poora instruction path execute karta hai (predication, branching nahi).

Worked example Warp divergence (the performance killer)

__global__ void divergentKernel(int *data) {
    int i = threadIdx.x;
    if (i % 2 == 0) {
        data[i] = data[i] * 2;      // Even threads
    } else {
        data[i] = data[i] + 100;    // Odd threads
    }
}

Kya hota hai:

Threads 0, 2, 4, ..., 30 if branch lete hain
Threads 1, 3, 5, ..., 31 else branch lete hain
Warp nahi kar sakta dono paths simultaneously execute (single instruction stream)
Serialization: Warp data[i] * 2 execute karta hai threads 1,3,5,... ko masked off (idle) karke, phir data[i] + 100 execute karta hai threads 0,2,4,... ko masked off karke
Effective throughput: Half (ya zyada branches hone par aur bhi bura)

Ye step kyun? Hardware predication implement karta hai: dono paths fetch hoti hain, lekin har thread sirf apne active path ke results commit karta hai. Isse is warp ke liye execution time double ho jaata hai.

Common mistakes

Common mistake "Zyada CUDA cores = zyada fast"

Kyun sahi lagta hai: Spec sheets core counts advertise karte hain (RTX 4090: 16,384 cores!). Zyada cores = zyada parallelism = zyada fast.

Kyun galat hai: Performance in chezon se limited hoti hai:

Memory bandwidth: Agar tumhara kernel memory-bound hai (data ka intezaar karne mein cycles spend kar raha hai), toh cores badhane se koi fayda nahi. Bottleneck DRAM speed hai (~1 TB/s), compute nahi (100 TFLOPS).
Warp divergence: Branchy code serialize ho jaata hai, cores waste hote hain.
Occupancy: Agar tumhara kernel bahut zyada registers use kare, toh SM latency hide karne ke liye enough warps fit nahi kar sakta.

Sahi intuition: Core count compute-bound workloads (dense matrix multiply, convolutions) ke liye ceiling set karta hai. Real applications ke liye, memory optimization (coalescing, tiling, shared memory) zyada matter karta hai.

Fix: nvprof ya Nsight Compute se profile karo. Agar "memory throughput" ~100% ke paas hai, toh memory patterns optimize karo, sirf cores mat badhao.

Recall Ek 12-saal ke bachche ko explain karo

Socho ek factory hai jahan tumhe 10,000 toy cars paint karni hain, har ek alag color mein car number ke hisaab se.

CPU approach: Tum 8 bahut smart painters hire karte ho. Har painter instruction sheet padh sakta hai, exact paint color mix kar sakta hai, ek car perfectly paint kar sakta hai, aur next par move kar sakta hai. Har car ke liye fast, lekin ek time par sirf 8.

GPU (CUDA) approach: Tum 16,384 simple painters hire karte ho, lekin ye "warps" ki teams of 32 mein kaam karte hain. Har warp ko ek time par EK instruction milti hai: "Sab apni car red paint karo!" Saare 32 painters simultaneously red paint karte hain. Next instruction: "Blue stripes add karo!" Saare 32 sath mein blue stripes add karte hain.

Problem: Kya hoga agar car #1 ko red chahiye aur car #2 ko blue? Warp split nahi ho sakta. Pehle, painter #1 red paint karta hai jabki #2 wait karta hai (kuch nahi karta). Phir #2 blue paint karta hai jabki #1 wait karta hai. Ye hai "warp divergence" — team ka aadha idle hai, toh tum effectively sirf 8 painters use kar rahe ho, 16 nahi.

Trick: Kaam organize karo taaki ek warp ke saare 32 painters ek hi kaam karein. Agar cars 0-31 sab same paint pattern paate hain (ya bahut similar), toh GPU CPU ko crush karta hai. Isliye GPUs graphics mein amazing hain (har pixel similar shading math karta hai) aur AI mein (har neuron same multiply-add karta hai) lekin messy list sort karne mein bure hain jahan har item ko custom logic chahiye.

Connections

6.2.01-GPU-vs-CPU-architecture: GPUs ne MIMD ki jagah SIMT kyun choose kiya
6.2.02-GPU-memory-hierarchy: Shared memory aur coalescing CUDA core efficiency ko kaise affect karte hain
6.2.04-occupancy-and-performance: Latency hide karne ke liye warp occupancy maximize karna
7.3.01-parallel-programming-patterns: Map/reduce patterns jo warp execution exploit karte hain
9.1.02-neural-network-training-on-GPUs: Matrix multiply CUDA ke liye perfectly suited kyun hai
6.1.03-SIMD-vs-SIMT: GPU SIMT ko CPU SIMD vectorization se compare karna

#flashcards/hardware

CUDA core kya hota hai aur ye CPU core se kaise alag hai?

Ek CUDA core ek general-purpose ALU hai (NVIDIA ka marketing term) jo ek clock cycle mein ek FP32 ya integer operation execute karta hai — sabse notable FMA hai. CPU core ke unlike, iske paas independent instruction decode, branch prediction, aur control logic nahi hoti — ye ek warp mein 32 cores mein shared hoti hain. Ek CPU core independent instruction streams run karta hai; CUDA cores apne warp ke andar lockstep mein execute hote hain.

Kya SM mein zyada CUDA cores add karne se ek single warp instruction ki latency kam hoti hai?

Nahi. Jab tak SM mein kam se kam 32 ALUs hain (warp mein har thread ke liye ek), ek single warp instruction ek cycle mein complete hoti hai. Extra cores throughput add karte hain (multiple schedulers ke zariye per cycle zyada warps issue hoti hain), ek warp ke liye latency kam nahi hoti.

CUDA execution model mein warp kya hota hai?

Warp 32 threads ka ek group hai jo ek hi instruction simultaneously execute karte hain. Ek warp ke threads control logic share karte hain aur lockstep mein execute hote hain (SIMT). Agar threads diverge karein (alag branch paths), toh warp un paths ko serialize karta hai, efficiency kam hoti hai.

1D CUDA grid ke liye global thread ID derive karo

globalID = blockIdx.x * blockDim.x + threadIdx.x. Ye ek thread ki apne block mein position (threadIdx) ko uske pehle aane wale complete blocks ki sankhya se offset karke (blockIdx * blockDim) ek unique global position par map karta hai.

Warp divergence ek performance problem kyun hai?

Jab ek warp ke threads alag execution paths lete hain (if/else), warp unhe serialize karta hai — ek path execute karta hai non-participating threads ko mask karke, phir doosra. N divergent paths ke liye ye effective throughput 1/N tak cut kar sakta hai, kyunki cores un paths ke dauran idle rehte hain jo wo nahi lete.

Memory latency fully hide karne ke liye zaruri warps ka formula kya hai, aur har symbol ka kya matlab hai?

T_{required} = \frac{S \cdot L}{C}

, jahan

S

= warp schedulers ki sankhya,

L

= memory stall cycles,

C

= stalls ke beech per warp useful compute cycles. Compute intensity

C

badhane se zaruri warps ki sankhya kam hoti hai.

Memory transaction efficiency kya hai, aur kya ye ek bandwidth hai?

η = \frac{useful bytes}{transactions \times 128}

hai, ek dimensionless fraction (bytes ÷ bytes), bandwidth nahi. Effective bandwidth =

η \times

peak DRAM bandwidth. Coalesced access →

η \approx 1

; strided access

η

ko ~3% tak gira sakta hai.

256 threads ke ek block mein kitne warps hote hain?

8 warps (256 ÷ 32).

Warp size 32 threads kyun hai?

Ye ek hardware convention hai: 2 ki power (masking/indexing simplify karta hai) jo 8/16/32 ki SIMD-width lineage se match karta hai. Ye kisi SM ke core count se DERIVED NAHI hai. Warp size register-file occupancy limits ke saath bhi interact karta hai.

SIMT kya hai aur ye SIMD se kaise alag hai?

SIMT (Single Instruction Multiple Thread): ek instruction multiple threads par operate karti hai, har ek ke apne registers hote hain, jo theoretically diverge kar sakte hain. SIMD: ek instruction ek single register mein packed data par, bina per-lane divergence ke. SIMT independent threads ka illusion deta hai jabki unhe lockstep mein execute karta hai.