Thread blocks and grids — Hinglish notes

6.2.6 · Hardware › GPU Architecture

The Two-Level Hierarchy

GPUs millions of threads simultaneously execute karte hain, lekin hardware constraints ek structure force karti hain:

Thread Block (Cooperative Thread Array):

threads ka ek group (typically 128–1024) jo same kernel code execute karta hai
Ek block ke saare threads same Streaming Multiprocessor (SM) par run karte hain
Fast on-chip shared memory share karte hain (user-managed cache)
__syncthreads() barrier use karke synchronize kar sakte hain
Maximum size: hardware-limited (e.g., modern GPUs par 1024 threads/block)

Grid:

thread blocks ka ek collection jo saath mein launch hota hai
Blocks 1D, 2D, ya 3D arrays ho sakte hain (problem geometry se match karte hue)
Blocks independently execute hote hain aur kisi bhi order mein run ho sakte hain
GPU ke saare SMs ko span karta hai
Ek kernel ke andar cross-block synchronization nahi hoti (blocks autonomous units hote hain)

Why This Design?

Blocks kyon exist karte hain:

Scalability: Same kernel 10 SMs wale GPUs par bhi run karta hai aur 100 SMs wale par bhi—runtime bas blocks ko differently schedule karta hai
Data locality: Nearby data elements process karne wale threads physically close rehte hain, fast memory share karte hue
Synchronization: Sirf wahi threads jo coordinate karne ki zaroorat hain (ek block ke andar) synchronize kar sakte hain—global stalls avoid hote hain

Grids kyon exist karti hain:

Problem decomposition: Data structures ke saath natural mapping (images ke liye 2D grid, volumes ke liye 3D)
Parallelism exposure: Programmer maximum parallelism express karta hai; hardware jo exploit kar sake karta hai
Load balancing: Runtime blocks ko SMs par dynamically distribute karta hai, idle time hide karta hai

Deriving Global Thread IDs

1D problem ke liye (e.g., length N ka vector addition):

Setup:

Grid mein G blocks hain, har ek mein B threads hain
Total threads: $T = G \times B$
Har thread ek element process karta hai

Yeh indexing kyon? Block 0 mein thread 0 element 0 se map hona chahiye, block 1 mein thread 0 element B se map hona chahiye, etc.

Derivation:

Block 0: threads [0, B-1]      → elements [0, B-1]
Block 1: threads [0, B-1]      → elements [B, 2B-1]
Block k: threads [0, B-1]      → elements [kB, (k+1)B-1]

Block k mein thread t ke liye: $global_id = k \cdot B + t = blockIdx.x \times blockDim.x + threadIdx.x$

Yeh step kyon? Block index humein batata hai ki pehle kitne complete blocks aa chuke hain (har ek B elements contribute karta hai), aur threadIdx current block ke andar hamari position deta hai.

Worked example Vector Addition (1D)

Do arrays add karo: C[i] = A[i] + B[i] for i = 0..N-1

// Host code: launch kernel
int N = 1000000;
int threadsPerBlock = 256;
int blocks = (N + threadsPerBlock - 1) / threadsPerBlock;  // ceiling division
vecAdd<<<blocks, threadsPerBlock>>(A, B, C, N);
// Kernel
__global__ void vecAdd(float *A, float *B, float *C, int N) {
   int i = blockIdx.x * blockDim.x + threadIdx.x;
   if (i < N) {  // Guard for oversubscription
       C[i] = A[i] + B[i];
   }
}

Ceiling division kyon? Agar N=1000 aur block size=256 hai, to saare elements cover karne ke liye 4 blocks chahiye (3 nahi). Formula: $⌈ N / B ⌉ = (N + B - 1) / B$ integer math.

If-guard kyon? Last block mein extra threads ho sakte hain (block 4 mein 256 threads hain lekin sirf 232 elements bache hain). Guard ke bina, threads invalid memory access karte hain.

Worked example Matrix Multiplication (2D)

Compute karo C = A × B jahan C[row][col] = sum(A[row][k] * B[k][col])

// Launch: 2D grid of2D blocks
dim3 threadsPerBlock(16, 16);  // 256 threads
dim3 blocks((N + 15)/16, (M + 15)/16);
matMul<<<blocks, threadsPerBlock>>>(A, B, C, M, N, K);
__global__ void matMul(float *A, float *B, float *C, int M, int N, int K) {
   int row = blockIdx.y * blockDim.y + threadIdx.y;
   int col = blockIdx.x * blockDim.x + threadIdx.x;
   if (row < M && col < N) {
       float sum = 0.0f;
       for (int k = 0; k < K; k++) {
           sum += A[row * K + k] * B[k * N + col];
       }
       C[row * N + col] = sum;
   }
}

2D blocks kyon? Har output element C[row][col] independent hai—threads ko outputs se directly map karna natural hai. 16×16 blocks sweet spot hit karte hain: latency hide karne ke liye enough threads, shared memory mein fit hone ke liye kaafi small (optimized versions ke liye).

Rows ke liye blockIdx.y kyon? Convention: x-axis horizontal hai (columns), y-axis vertical hai (rows mathematical matrix notation se match karte hue).

Block Scheduling

Warp-based execution within blocks:

Blocks ke andar threads warps (32 threads) mein grouped hote hain
Ek warp ke saare threads same instruction simultaneously execute karte hain (SIMT model)
Efficiency ke liye block size 32 ka multiple hona chahiye

Example: 256-thread block = 8 warps

Agar threads diverge ho jaayein (if-else), kuch warp threads mask ho jaate hain—phir bhi cycles consume karte hain
__syncthreads() ke zariye synchronization ensure karti hai ki block ke saare warps barrier tak pahunche kisi ke aage badhne se pehle

Resource limits per SM:

Maximum threads (e.g., Ampere par 2048)
Registers per thread
Shared memory per block
Maximum blocks (typically 16-32)

Block size kyon matter karta hai:

Too small (e.g., 32): SM underutilize hota hai, memory latency hide nahi hoti
Too large (e.g., 1024): SM par fewer blocks fit hote hain, occupancy reduce hoti hai agar ek block stall kare
Sweet spot: 128–512 threads, empirically ya occupancy calculator se choose karo

Common mistake "Blocks Synchronize Automatically"

Galat idea: Yeh assume karna ki block 0 block 1 se pehle finish ho jaata hai, ya global memory use karke same kernel ke andar blocks ke beech "messages pass" karo.

Kyon sahi lagta hai: Blocks sequentially numbered hain (blockIdx.x = 0, 1, 2...), jo order suggest karta hai.

Fix yeh hai: Blocks arbitrary order mein run hote hain aur time mein overlap kar sakte hain. Hardware block 5 ko block 1 se pehle schedule kar sakta hai agar SM2 idle ho. Agar tumhe cross-block coordination chahiye:

Option 1: Multiple kernel launches mein split karo (implicit global barrier)
Option 2: Grid-level synchronization use karo (CUDA 9+, cooperative groups)
Option 3: Global memory par atomic operations (expensive)

Example: Yeh mat karo:

// Block 0 writes flag=1, other blocks wait—DEADLOCK!
if (blockIdx.x == 0) {
   work();
   flag[0] = 1;
} else {
   while (flag[0] == 0);  // ❌ May run before block 0!
}

Recall Explain to a 12-Year-Old

Imagine karo tum ek giant school art project organize kar rahe ho: 10,000 students ek massive mural paint kar rahe hain. Tum sab ko wall par swarm nahi karne de sakte—chaos ho jaayega! Isliye:

Students ko 30 ke teams mein split karo (thread blocks). Har team ko wall ka ek section milta hai aur apni table par shared paint buckets milte hain (shared memory).
Teams ek doosre se baat kar sakte hain ("Hey, kya sab log sky ke saath done hain?") lekin sirf apni table ke andar—gym ke across cheekh ke bolna bahut slow hai.
Gym mein 10 workstations hain (SMs), isliye 10 teams simultaneously kaam karti hain. Jaise hi teams finish karti hain, teacher un workstations ko naye sections assign karta hai.
Grid poora mural hai jo sections mein divided hai, har team ke liye ek.

GPUs isi tarah kaam karte hain: millions of tiny workers (threads) cooperative teams (blocks) mein grouped hote hain jo canvas (grid) ke alag-alag parts paint karte hain bina door wali teams ke saath coordinate karne ki zaroorat ke. Isi tarah tumhara GPU ek game frame billions of pixels ke saath milliseconds mein render karta hai!

Connections

CUDA Programming Model – grids/blocks launch karne ke liye syntax aur API
Streaming Multiprocessors – hardware jo thread blocks execute karta hai
Warp Execution – ek block ke andar threads SIMT groups mein kaise run karte hain
Shared Memory – intra-block communication ke liye fast on-chip storage
Thread Synchronization – __syncthreads() aur cooperative groups
GPU Occupancy – SM par active threads maximize karna
Memory Coalescing – thread indexing memory bandwidth ko kaise affect karta hai
Parallel Algorithm Design – problems ko blocks ke grids mein decompose karna

#flashcards/hardware

GPU execution hierarchy ke do main levels kya hain? :: Thread blocks aur grids. Threads blocks mein grouped hote hain jo resources share karte hain, aur blocks ek grid mein organized hote hain jo problem space ko cover karta hai.

Alag blocks ke threads ek kernel ke andar synchronize kyon nahi kar sakte?

Blocks independently execute hote hain aur kisi bhi SM par kisi bhi order mein run ho sakte hain. Yeh independence alag SM counts wale GPUs par scalability enable karti hai.

1D mein thread ke global ID ka formula kya hai?

globalID = blockIdx.x * blockDim.x + threadIdx.x. Yeh pichle saare blocks ke threads ko account karta hai plus current block mein position.

Grid size calculate karte waqt ceiling division kyon use karte hain?

Yeh ensure karne ke liye ki enough blocks saare data elements cover karein. Formula (N + B - 1) / B round up karta hai, isliye last block remaining elements handle karta hai chahe N block size se divisible na ho.

Thread identification ke liye built-in CUDA variables kya hain?

threadIdx (block mein position), blockIdx (grid mein block position), blockDim (block size), gridDim (grid size). Har ek ke x, y, z components hote hain.

32-multiple block size kyon recommend ki jaati hai?

Threads 32 ke warps mein execute hote hain. Non-multiples warp slots waste karte hain aur SIMT execution model ki wajah se efficiency reduce hoti hai.

SM par kitne blocks simultaneously run kar sakte hain yeh kya limit karta hai?

SM par total threads, registers per thread, shared memory per block, aur maximum concurrent blocks ki hardware limits. Sabse tight constraint occupancy determine karta hai.

Yeh assume kyon nahi kar sakte ki block 0 block 1 se pehle execute hota hai?

GPU runtime load balancing aur scalability ke liye blocks ko arbitrary order mein kisi bhi available SM par independently schedule karta hai.

Jab N block size se divisible nahi hota to extra threads ka kya hota hai? :: Last block mein excess threads hote hain. Out-of-bounds memory access prevent karne ke liye if-guard if (i < N) use karo.

Matrix operations ke liye 2D blocks kyon use karte hain?

Natural mapping: har thread (row, col) par ek output element compute karta hai. Problem geometry se match karta hai aur indexing logic simplify karta hai.

GPU occupancy kya hai?

SM par active warps ka maximum possible warps se ratio. Zyada occupancy parallelism ke through memory latency hide karne mein help karta hai.

Thread blocks ko "cooperative thread arrays" kyon kaha jaata hai?

Ek block ke threads shared memory aur synchronization barriers (__syncthreads()) ke through cooperate kar sakte hain, unlike alag blocks ke threads.