CUDA programming model basics — Hinglish notes

6.2.13 · Hardware › GPU Architecture

CUDA kya hai?

CUDA (Compute Unified Device Architecture) NVIDIA ka parallel computing platform hai jo programmers ko GPU ki power general-purpose computation ke liye use karne deta hai, sirf graphics ke liye nahi.

YEH kyun exist karta hai: Graphics cards lakhon pixels ko parallel mein render karne ke liye evolve hui. NVIDIA ne realize kiya ki yeh massive parallelism scientific, AI, aur data problems ko bhi solve kar sakta hai, agar woh programmers ko direct access de dein.

YEH kya provide karta hai: Ek C/C++ jaisi language extension + runtime libraries + compiler (nvcc) jo tumhe GPU ke liye code likhne deta hai.

YEH kaise kaam karta hai: Tum special functions (kernels) likhte ho jo GPU par run hote hain. Tum kaam ko threads, blocks, aur grids mein organize karte ho. GPU scheduler inhe apne hardware par distribute karta hai.

CUDA Execution Model: First Principles Se

Parallelism mein Structure kyun Zaroori hai

Fundamental problem: Agar tumhare paas 10,000 tasks aur 10,000 workers hain, tumhe chahiye:

Ek tarika identify karne ka ki kaun sa worker kaun sa task karta hai
Ek tarika coordinate karne ka jab workers ko information share karni ho
Ek tarika map karne ka logical work organization ko physical hardware par

CUDA ise teen-level hierarchy se solve karta hai:

Formula CUDA Hierarchy (Logical → Physical)

Logical Organization: $Grid \to Blocks \to Threads$

Physical Mapping: $Grid maps to GPU \to Blocks map to SMs \to Threads map to CUDA cores$

Teen levels kyun?

Thread level: Individual work items ko identify karta hai
Block level: Related threads mein cooperation enable karta hai (shared memory, barriers)
Grid level: Puri GPU par hardware assumptions ke bina scale karta hai

DERIVATION - Thread Indexing:

1D problem ke liye jisme total N elements hain:

Launch configuration: <<numBlocks, threadsPerBlock>>>
Har thread apna global index compute karta hai: $globalIdx = blockIdx.x \times blockDim.x + threadIdx.x$

Yeh formula kyun?

blockIdx.x: Main kaun se block mein hoon? (0, 1, 2, ...)
blockDim.x: Ek block mein kitne threads hain? (jaise 256)
blockIdx.x × blockDim.x: Mere block se pehle kitne threads aa chuke hain?
threadIdx.x: Mere block ke andar meri position (0 se blockDim.x-1 tak)
Sum = poore grid mein meri unique position

2D problems (images) ke liye: $row = blockIdx.y \times blockDim.y + threadIdx.y$ $col = blockIdx.x \times blockDim.x + threadIdx.x$ $globalIdx = row \times width + col$

Memory Hierarchy

Memory kyun matter karta hai: Data move karna expensive hai. Ek memory access math operation se 100× zyada time le sakta hai. CUDA multiple memory types provide karta hai jisme alag-alag speed/scope tradeoffs hain.

Formula CUDA Memory Hierarchy (Fast → Slow)

| Memory Type | Scope | Lifetime | Speed | Size | |-------------|----------|----| | Registers | Per-thread | Thread lifetime | 1 cycle | ~64KB/SM | | Shared Memory | Per-block | Block lifetime | ~5 cycles | 48-164KB/SM | | L1/L2 Cache | Per-SM / GPU | Automatic | ~10-30 cycles | 128KB-40MB | | Global Memory | All threads | Program lifetime | ~200-400 cycles | GBs (VRAM) | | Host Memory | CPU | Persistent | ~500-1000 cycles | System RAM |

Speed reasoning first principles se:

Registers: On-chip, seedha compute units mein → wire distance ~mm → picoseconds
Shared memory: On-chip SRAM, SM ke andar shared → phir bhi mm → few ns
Global memory: Off-chip GDDR/HBM → wire distance ~cm + memory controller → 100+ ns
Host memory: PCIe bus ke paar → meters of signaling → microseconds

Access pattern formula: $Effective Bandwidth = \frac{Theoretical BW \times Coalescing Efficiency}{Bank Conflicts + Unaligned Access Penalties}$

Apna Pehla CUDA Kernel Likhna

Vector Addition: CUDA ka "Hello World"

Problem: Do arrays add karo: C[i] = A[i] + B[i] i = 0 se N-1 tak

CPU version (sequential):

for (int i = 0; i < N; i++) {
    C[i] = A[i] + B[i];
}
// Time: O(N) - har iteration mein ek element

CUDA version (parallel):

Worked example Complete CUDA Vector Addition

// KERNEL: GPU device par run hota hai
__global__ void vectorAdd(float *A, float *B, float *C, int N) {
   // Har thread apna unique index compute karta hai
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   
   // Boundary check: sirf valid indices process karo
   if (idx < N) {
       C[idx] = A[idx] + B[idx];
   }
}
int main() {
   int N = 1000000;
   size_t bytes = N * sizeof(float);
   
   // 1. HOST MEMORY ALLOCATE KARO
   float *h_A = (float*)malloc(bytes);
   float *h_B = (float*)malloc(bytes);
   float *h_C = (float*)malloc(bytes);
   
   // h_A, h_B ko data se initialize karo...
   
   // 2. DEVICE MEMORY ALLOCATE KARO
   float *d_A, *d_B, *d_C;
   cudaMalloc(&d_A, bytes);
   cudaMalloc(&d_B, bytes);
   cudaMalloc(&d_C, bytes);
   
   // 3. HOST → DEVICE COPY KARO
   cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
   cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);
   // 4. KERNEL LAUNCH KARO
   int threadsPerBlock = 256;
   int numBlocks = (N + threadsPerBlock - 1) / threadsPerBlock;
   vectorAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, N);
   
   // 5. DEVICE → HOST COPY KARO
   cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
   
   // 6. CLEANUP
   cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
   free(h_A); free(h_B); free(h_C);
}

Har step kyun? Step 1 (Host allocation): Input/output data ke liye CPU memory Step 2 (Device allocation): GPU ka memory space alag hota hai—VRAM allocate karna zaroori hai

Step 3 (Copy H→D): GPU system RAM access nahi kar sakta—PCIe ke zariye data copy karo

Step 4 (Kernel launch):

<<numBlocks, threadsPerBlock>>>: Execution configuration
numBlocks = ceil(N / 256): Saare N elements cover karne ke liye enough threads chahiye
256 threads/block kyun? Hardware constraint: blocks mein max threads hote hain (1024). 256 ek accha balance hai: SM par 4-8 blocks allow karta hai, warps (32) evenly fill karta hai numBlocks ka formula: $numBlocks = ⌈ \frac{N}{threadsPerBlock} ⌉ = \frac{N + threadsPerBlock - 1}{threadsPerBlock}$ Ceiling division trick floating-point math se bachata hai

Step 5 (Copy D→H): Results CPU memory mein wapas lao

Step 6 (Cleanup): Host aur device dono ki memory free karo

Boundary Check: `if (idx < N)` kyun?

Problem: Agar N = 1000 aur threadsPerBlock = 256 hai:

4 blocks × 256 threads = 1024 threads total chahiye
Lekin hamare paas sirf 1000 elements hain
Aakhri 24 threads (idx 1000-1023) extra hain

Check ke bina: Yeh threads A[1000], A[1001], ... access karte → out of bounds! → crash ya garbage data

Check ke saath: Extra threads kuch nahi karte (idle rehte hain). Thoda performance cost, disaster se bachata hai.

Common mistake Common CUDA Mistakes

Mistake 1: Boundary check bhool jaana

__global__ void bad(float *A, int N) {
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   A[idx] = idx;  // CRASH agar idx >= N!
}

Kyun sahi lagta hai: "Maine bilkul enough threads launch kiye, hai na?" FIX: Tumne blocks launch kiye, aur blocks fixed sizes mein aate hain. Hamesha if (idx < N) check karo.

Mistake 2: Galat memory copy direction

cudaMemcpy(d_A, h_A, bytes, cudaMemcpyDeviceToHost);  // GALAT!

Kyun sahi lagta hai: "d_A destination hai, h_A source hai..." FIX: Enum data flow ki direction specify karta hai: DeviceToHost matlab GPU → CPU. Input ke liye HostToDevice use karo.

Mistake 3: CUDA errors check na karna

cudaMalloc(&d_A, bytes);  // Fail ho sakta hai! Koi check nahi!

Kyun sahi lagta hai: "Yeh toh sirf allocation hai, kaam kar jayega." FIX: GPU memory khatam ho sakti hai. Hamesha check karo:

cudaError_t err = cudaMalloc(&d_A, bytes);
if (err != cudaSuccess) {
   printf("Error: %s\n", cudaGetErrorString(err));
}

Mistake 4: 2D indexing ke liye galat formula use karna

int idx = blockIdx.x * blockDim.y + threadIdx.x;  // GALAT!

Kyun sahi lagta hai: "Main threads ke liye x use kar raha hoon..." FIX: Block dimensions .x, .y, .z use karte hain lekin block index dimensions se match karne chahiye. 1D grid ke liye: puri jagah .x use karo. 2D ke liye: rows ke liye blockIdx.y × blockDim.y + threadIdx.y.

Synchronization aur Cooperation

Block ke Andar Thread Synchronization

Kyun zaroori hai: Ek block ke threads resources share karte hain (shared memory). Synchronization ke bina, race conditions hoti hain.

Worked example Shared Memory Use Karna (Reduction)

Problem: Ek array ko parallel mein sum karo

__global__ void sumReduce(float *input, float *output, int N) {
   __shared__ float shared[256];  // Is block ke liye shared memory
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   int tid = threadIdx.x;
   
   // Data ko shared memory mein load karo
   shared[tid] = (idx < N) ? input[idx] : 0.0f;
   __syncthreads();  // Saare threads ke load hone ka wait karo
   
   // Shared memory mein reduction
   for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
       if (tid < stride) {
           shared[tid] += shared[tid + stride];
       }
       __syncthreads();  // Agli iteration se pehle wait karo
   }
   // Thread 0 result likhta hai
   if (tid == 0) {
       output[blockIdx.x] = shared[0];
   }
}

Yeh kyun kaam karta hai:

__shared__ declaration: Fast on-chip memory allocate karta hai jo block ke saare threads ko visible hai
Pehla __syncthreads(): Ensure karta hai ki saare threads loading finish kar lein pehle koi kuch padhne se
Reduction loop: Har iteration mein active threads ko half karta hai

Stride 128: threads 0-127 pairs add karte hain
Stride 64: threads 0-63 pairs add karte hain
...jab tak sirf thread 0 ke paas sum na ho

Doosra __syncthreads(): Ensure karta hai ki pichli iteration ke writes agli iteration ke reads se pehle visible ho jayein
Thread 0 likhta hai: Sirf ek thread block ka result global memory mein likhta hai

Derivation - Reduction O(log N) kyun hai: $Iterations = lo g_{2} (blockDim.x)$ Har iteration kaam ko half karta hai. 256 threads ke liye: 255 sequential additions ki jagah 8 iterations.

Global Synchronization

Key constraint: __syncthreads() sirf ek block ke andar kaam karta hai. Alag-alag blocks ke threads kernel execution ke dauran synchronize nahi kar sakte.

Kyun: Blocks kisi bhi order mein execute ho sakte hain, ya alag SMs par concurrently bhi. Inter-block synchronization allow karna deadlocks create karta.

Blocks ke paar synchronize kaise karein: Alag kernels launch karo. GPU guarantee karta hai ki kernel 1 ke saare threads kernel 2 ke kisi bhi thread se pehle finish ho jayenge.

Performance Considerations

Memory Coalescing

Coalesced access: Ek warp (32 consecutive threads) ke threads consecutive memory addresses access karte hain → GPU ek transaction mein combine karta hai.

Uncoalesced access: Threads scattered addresses access karte hain → GPU kai choti transactions karta hai → 10-100× slower.

Formula Coalescing Rules

Perfect coalescing (32 threads ke liye ek 128-byte transaction):

float *data;
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float value = data[idx];  // Sequential access

Terrible performance (32 separate transactions):

float value = data[idx * 1000];  // 1000 ka stride

Kyun: GPU memory controllers aligned 32/64/128-byte chunks par kaam karte hain. Agar thread 0 byte 0 padhta hai, thread 1 byte 4, thread 2 byte 8, ..., thread 31 byte 124 → sab ek 128-byte chunk mein → ek transaction.

Agar accesses scattered hain, har thread ki request alag chunk hit karta hai → 32 transactions.

Occupancy

Occupancy active warps aur SM par maximum possible warps ka ratio hai.

Kyun matter karta hai: Zyada occupancy → zyada warps available → better latency hiding.

Occupancy limit karne wale factors:

Registers per thread: Zyada registers → kam threads fit hote hain
Shared memory per block: Zyada shared memory → kam blocks fit hote hain
Block size: Bahut chota → resources waste; bahut bada → SM par blocks limit karta hai

$Occupancy = \frac{Active Warps per SM}{Maximum Warps per SM}$

Typical goal: 50-75% occupancy aksar enough hota hai. 100% hamesha acchi performance ke liye zaroori nahi hai.

Recall Feynman: 12-saal ke bachche ko samjhao

Socho tumhare paas 10,000 problems wali ek badi math worksheet hai. Agar tum khud karo, toh bahut time lagega—ek problem ek time par.

Ab socho tumhare paas 10,000 students wala ek classroom hai, aur tum har student ko exactly ek problem dete ho. Sab ek saath kaam karte hain, aur achanak poori worksheet utne hi time mein ho jaati hai jitna ek problem solve karne mein lagta hai!

Yahi GPU karta hai. "Students" threads hain—chote workers. CUDA woh tarika hai jisse tum unhe batate ho ki kaun sa problem solve karna hai.

Tum students ko groups (blocks) mein organize karte ho jo tables par baithe hain. Ek hi table par baithe students ek whiteboard (shared memory) share kar sakte hain milkar kaam karne ke liye. Lekin alag tables par baithe students tab tak baat nahi kar sakte jab tak sab done na ho jayein.

Tumhe dhyan rakhna hoga: agar tumhare paas 10,000 problems hain lekin 10,240 students hain, toh aakhri 240 ke paas solve karne ke liye koi problem nahi hai. Tumhe unhe bolna hoga "is baar bahar baitho" warna woh imaginary problems solve karne ki koshish karenge aur sab garbad kar denge!

Connections

GPU Architecture Overview - CUDA ke liye hardware foundation
Streaming Multiprocessors - Physical execution units jo blocks run karte hain
Memory Hierarchy - Detailed memory performance analysis
Thread Warps and SIMT - Threads hardware par actually kaise execute hote hain
Parallel Algorithm Design - Problems decompose karne ki strategies
OpenCL vs CUDA - Alternative programming models
Deep Learning Frameworks - TensorFlow/PyTorch hood ke neeche CUDA use karte hain

Flashcards

#flashcards/hardware

CUDA kernel kya hota hai? :: Ek function jo GPU par run hota hai, bahut saare threads dwara parallel mein execute kiya jaata hai. __global__ keyword se declare kiya jaata hai.

1D mein global thread index ka formula kya hai?

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

CUDA kernels mein if (idx < N) kyun zaroori hai? :: Kyunki hum fixed size ke blocks launch karte hain (jaise 256 threads), hamare paas data elements se zyada threads ho sakte hain. Extra threads ko out-of-bounds memory access karne se rokna zaroori hai.

CUDA mein host aur device mein kya fark hai?

Host = CPU aur system RAM. Device = GPU aur VRAM. Inke alag memory spaces hain jinhein explicit copies ki zaroorat hoti hai.

<<numBlocks, threadsPerBlock>>> kya specify karta hai?

Kernel launch configuration: kitne blocks launch karne hain aur har block mein kitne threads.

__syncthreads() kis kaam aata hai?

Threads ko ek block ke andar synchronize karne ke liye. Ensure karta hai ki saare threads barrier tak pahunch jayein pehle koi aage badhe. Tab use hota hai jab threads shared memory ke zariye data share karte hain.

Kya alag-alag blocks ke threads kernel execution ke dauran synchronize kar sakte hain?

Nahi. Blocks independently execute hote hain aur kisi bhi order mein run ho sakte hain. Blocks ke paar synchronize karne ka ek hi tarika hai: alag kernels launch karo.

Memory coalescing kya hai?

Jab ek warp ke threads consecutive memory addresses access karte hain, GPU requests ko ek transaction mein combine karta hai. Non-coalesced access kai alag transactions cause karta hai aur bahut slow hota hai.

__shared__ memory kya hai?

Ek block ke saare threads dwara shared fast on-chip memory. Global memory se bahut faster (~5 cycles vs ~400 cycles) lekin size mein limited (48-164KB per SM).

CUDA mein warp kya hota hai?

32 threads ka ek group jo lockstep mein execute hota hai (SIMT). NVIDIA GPUs par fundamental execution unit.

N elements ke liye blocks ki zaroorat kaisi calculate karte hain?

numBlocks = (N + threadsPerBlock - 1) / threadsPerBlock (integer arithmetic use karke ceiling division)

CUDA ke execution hierarchy ke teen levels kya hain?

Grid → Blocks → Threads. Grid GPU par map hota hai, blocks SMs par map hote hain, threads CUDA cores par map hote hain.

threadsPerBlock ke liye 256 common choice kyun hai?

Yeh 32 (warp size) ka multiple hai, SM par 4-8 blocks allow karta hai, aur 1024 max thread limit se kaafi neeche rehta hai. Occupancy ke liye achha balance hai.

GPU occupancy kya hota hai?

Active warps aur SM par maximum possible warps ka ratio. Zyada occupancy warp switching ke zariye better latency hiding enable karta hai.

CUDA mein sabse fast memory type kaun si hai?

Registers (per-thread, ~1 cycle access), phir shared memory (~5 cycles), phir L1/L2 cache, phir global memory (~400 cycles).