CUDA programming model basics: Question bank — Hinglish deep dive

6.2.13 · D5 · Hardware › GPU Architecture › CUDA programming model basics

Yahan use hone wala har term (host, device, kernel, thread, block, grid, warp, coalescing, boundary check) parent note mein build kiya gaya hai — agar koi word unfamiliar lage, pehle wahan define karo. Related depth: Thread Warps and SIMT, Memory Hierarchy, Streaming Multiprocessors.

Pictures to lean on

Neeche teen traps geometric hain — jawab dete waqt inhe open rakho.

Teen-level hierarchy aur 2D global index kaise build hota hai (yeh "2D indexing" traps mein use hota hai):

Warps ek SM par kaise interleave hote hain — latency-hiding timeline (yeh "warp scheduling" traps mein use hota hai):

Pura memory-space map — registers, shared, constant, texture, local, global (yeh "which memory" traps mein use hota hai):

True or false — justify karo

Ek kernel launch code ko ek baar run karta hai, aur GPU usse internally loop karta hai.

False. Ek launch har index ke liye ek thread spawn karta hai jo sab kernel body simultaneously run karte hain; koi internal loop nahi hota — tumne loop ko hazaron parallel copies se replace kar diya.

Grid mein har thread shared memory ke through data share kar sakta hai.

False. Shared memory per-block hoti hai; sirf same block ke andar ke threads usse dekh sakte hain. Cross-block sharing slow global memory se karni padti hai.

threadIdx.x pure grid mein unique hota hai.

False. Yeh sirf ek block ke andar unique hota hai (0 se blockDim.x-1 tak); har block ise 0 se restart karta hai. Globally unique id hai blockIdx.x * blockDim.x + threadIdx.x.

2D grid ke liye, threadIdx.x akela image mein meri row batata hai.

False. 2D mein tumhe dono axes chahiye: row = blockIdx.y*blockDim.y + threadIdx.y aur col = blockIdx.x*blockDim.x + threadIdx.x — figure s01 ka left panel dekho.

Ek 2D block ek genuinely alag hardware cheez hai 1D block se.

False. dim3 block/grid shapes purely ek labelling convenience hain; hardware unhe phir bhi warps of 32 mein flatten karta hai. 2D/3D sirf image aur volume indexing ko readable banata hai.

256 threads wala block hamesha literally ek hi instant mein sare 256 run karta hai.

False. Threads warps of 32 mein execute hote hain; ek 256-thread block 8 warps hai jo SM schedule karta hai, possibly interleaved (figure s02), ek giant simultaneous step nahi.

Warps interleave karna GPU ko slower banata hai kyunki yeh keep switching karta hai.

False. Switching SM par free hai aur exactly isi tarah yeh memory latency hide karta hai: jab ek warp global memory ke liye ~300 cycles wait karta hai, doosra warp compute karta hai — figure s02 dikhata hai stalls kaise fill hote hain.

Agar main elements se kam threads launch karun, toh baaki elements silently aur safely skip ho jaate hain.

True par dangerous. Woh skip hote hain, lekin "safely" sirf tab agar tumne intentionally skip kiya ho; usually iska matlab hai tumne under-launch kiya aur tumhari output array mein stale/garbage entries hain — jab tak tum grid-stride loop use nahi karte (neeche dekho).

cudaMalloc tumhe ek pointer deta hai jise CPU directly dereference kar sakta hai.

False. Yeh VRAM mein ek device pointer return karta hai; host par ise dereference karna undefined behaviour hai — data ko cudaMemcpy ke zariye PCIe bus cross karni padti hai pehle.

Registers CUDA mein sabse chhoti aur sabse slow memory hain.

False. Registers sabse fast hain (~1 cycle) — scope mein sabse chhote (per-thread), lekin speed aur scope alag axes hain. Sabse slow reachable memory PCIe ke across host RAM hai.

Constant memory har access pattern ke liye global memory se faster hai.

False. Constant memory sirf tab fast hoti hai jab ek warp ke sare threads same address padhte hain (yeh ek small cache se broadcast karta hai); divergent addresses serialize ho jaate hain aur benefit kho dete hain — figure s03.

"Local memory" thread ke paas register ki tarah rehti hai.

False. Naam ke bawajood, local memory off-chip global memory hai jo register spills aur thread-private arrays ke liye use hoti hai; yeh global jaisi slow hai. Naam scope refer karta hai, location nahi — figure s03.

Threads per block zyada use karna hamesha faster hota hai kyunki zyada parallelism.

False. Blocks 1024 threads par cap hote hain aur har block SM ke registers/shared memory ke liye compete karta hai; bahut bade blocks reduce karte hain kitne blocks per SM fit hote hain, occupancy hurt hoti hai.

Spot the error

int idx = threadIdx.x; multi-block launch mein global index ke roop mein use kiya.

Wrong — yeh blockIdx.x ko ignore karta hai. Har block indices 0..blockDim-1 likhega, toh sare blocks same elements stomp karenge aur array ka bada hissa kabhi touch nahi hoga.

2D image kernel ke liye, int idx = blockIdx.x*blockDim.x + threadIdx.x; pixel index ke roop mein use kiya.

Wrong — yeh sirf ek row ke columns cover karta hai. Flattened pixel index hai row*width + col, dono x aur y global indices use karke (figure s01, right panel).

cudaMemcpy(h_C, d_C, bytes, cudaMemcpyHostToDevice); kernel ke baad.

Wrong direction. Results wapas copy karna GPU→CPU hai, toh yeh cudaMemcpyDeviceToHost hona chahiye; jaise likha hai yeh GPU result ko host garbage se overwrite karta hai.

numBlocks = N / threadsPerBlock; N = 1000, threadsPerBlock = 256 ke liye.

Wrong — integer division 3 blocks = 768 threads deta hai, toh elements 768–999 kabhi compute nahi hote. Ceiling division use karo (N + tpb - 1) / tpb.

Kernel body: A[idx] = idx; bina if (idx < N) ke.

Wrong — last (over-launched) threads A ke end ke baad likhte hain, out-of-bounds memory corruption ya crash cause karte hain.

Ek grid-stride loop likha for (int i = idx; i < N; i++) ek bade array cover karne ke liye.

Wrong — har thread end tak walk karega aur same tail reprocess karega. Stride total number of threads honi chahiye: for (int i = idx; i < N; i += blockDim.x*gridDim.x).

vectorAdd<<<numBlocks, threadsPerBlock>>>(h_A, h_B, h_C, N); host pointers pass karte hue.

Wrong — kernel device par run hota hai aur sirf device memory touch kar sakta hai; tumhe d_A, d_B, d_C pass karna hai, host arrays nahi.

Device memory ko free(d_A); se free karna.

Wrong — free host allocations ke liye hai jo malloc se aaye hain; cudaMalloc se device memory ko cudaFree se release karna hoga.

Why questions

CUDA work ko blocks mein kyun split karta hai instead of ek flat pool of threads ke?

Kyunki blocks SMs se map hote hain aur ek cooperation boundary dete hain (shared memory + barriers) jabki same code alag SM counts wale GPUs par scale karne deta hai — hardware choose karta hai kitne blocks ek saath run hote hain.

Grid-stride loop kyun use karein jab ek plain one-thread-per-element launch already kaam karta hai?

Yeh kernel ko N se decouple karta hai: ek fixed grid of, maano, 4096 threads kisi bhi size ka array process kar sakta hai kyunki har thread blockDim.x*gridDim.x aage hop karta hai, toh same launch un arrays ko bhi handle karta hai jo GPU ke thread count se bade hain.

Warps ek SM par interleave kyun hote hain instead of ek ko khatam karke doosra shuru karne ke?

Latency hide karne ke liye (figure s02): global memory par ~300 cycles stalled warp compute units ko idle kar deta, toh scheduler ek ready warp swap in karta hai — bahut se resident warps ALUs ko busy rakhte hain.

Constant aur texture memory offer kyun karein jab global memory sab kuch hold kar sakti hai?

Yeh specialized caches hain: constant memory ek value puri warp ko cheaply broadcast karta hai, aur texture memory hardware interpolation ke saath 2D/3D spatially-local reads ke liye optimize hai — dono global memory se behtar hain apne pattern ke liye (figure s03).

Boundary check if (idx < N) kyun chahiye agar numBlocks correctly compute kiya?

Ceiling division over-launch karta hai: numBlocks * threadsPerBlock usually N se bada hota hai, toh kuch extra threads exist karte hain aur unhe kuch nahi karne ki instruction deni hoti hai.

Data device par copy kyun karein — GPU system RAM kyun nahi padh sakta?

GPU ka ek physically alag memory space hai (VRAM); host RAM tak pahunchna slow PCIe bus cross karta hai, isliye CUDA ek explicit cudaMemcpy require karta hai data ko fast device memory mein stage karne ke liye.

256 threads per block prefer kyun karein, maano 30 ke bajaye?

256 warp size 32 ka multiple hai, toh koi warp half-empty nahi run karta; 30 ka block apne single partial warp mein 32 mein se 2 lanes waste karta hai har step par.

Data move karna math karne se itna zyada costly kyun hota hai?

Global memory off-chip rehti hai (~cm of wire + ek memory controller), toh ek access ~200–400 cycles hoti hai versus on-chip arithmetic op ke liye ~1 — poori memory hierarchy is gap ko hide karne ke liye exist karti hai.

Thread cooperation sirf ek block ke andar kyun hoti hai, blocks ke across nahi?

Alag blocks alag times par alag SMs par run ho sakte hain (ya bilkul concurrently nahi), toh koi guaranteed shared clock ya memory nahi hai synchronize karne ke liye — sirf intra-block barriers safe hain.

Edge cases

Kya hota hai jab N, threadsPerBlock se chhota ho (jaise N = 10, tpb = 256)?

numBlocks 1 ho jaata hai, 256 threads launch hote hain; boundary check sirf threads 0–9 ko kaam karne deta hai aur baaki 246 idle baithe hain — correct, bas underutilized.

Kya ho agar N, threadsPerBlock se exactly divisible ho?

Koi threads over-launch nahi hote, toh boundary check kabhi fire nahi karta — lekin tum ise phir bhi rakhte ho, kyunki perfect divisibility par rely karna usi second mein toot jaata hai jab N change hota hai.

Kya ho agar N = 0 ho?

Ceiling division numBlocks = 0 deta hai, toh kernel ek empty grid ke saath launch hota hai aur simply kuch nahi karta — ek valid no-op, error nahi.

Kya ho agar ek grid-stride loop N se zyada threads ke saath launch ho?

Jinke starting idx >= N hain woh threads kabhi loop body mein enter nahi karte — i < N condition built-in boundary check hai, toh grid-stride loops kisi bhi launch size par safe hain.

Kya ho agar ek 2D block 32×32 = 1024 threads ho?

Yeh exactly per-block cap hai; legal hai, lekin ek aur thread ke liye bhi koi room nahi bacha, aur register/shared-memory pressure phir bhi per SM aise kam blocks force kar sakta hai.

Kya ho agar do threads coordination ke bina same global address par likhte hain?

Tumhe ek race condition milti hai; final value woh hai jo last mein land karta hai (undefined), isliye aise patterns ko atomics ya ek redesigned access map chahiye — Parallel Algorithm Design dekho.

Kya ho agar ek warp ke 32 threads ek if ke alag branches lein?

Warp dono paths sequence mein execute karta hai inactive lanes masked off karke — yeh branch divergence branches ko serialize karta hai aur performance cost karta hai (detail Thread Warps and SIMT mein).

Agar main cudaMalloc ka return code check karna bhool jaun aur woh fail ho, toh kya dikhta hai?

d_A ek invalid pointer rehta hai aur kernel baad mein silently memory corrupt karta hai ya crash karta hai — failure apne cause se bahut door surface karta hai, isliye har CUDA call error-check honi chahiye.

Recall Jaane se pehle self-test karo

Global unique thread id formula (1D)? ::: blockIdx.x * blockDim.x + threadIdx.x. width-wide image mein 2D pixel index? ::: (blockIdx.y*blockDim.y+threadIdx.y)*width + (blockIdx.x*blockDim.x+threadIdx.x). Grid-stride loop stride value? ::: blockDim.x * gridDim.x — launch kiye gaye threads ki total number. Konsa copy direction inputs GPU par load karta hai? ::: cudaMemcpyHostToDevice. Boundary check exist kyun karta hai? ::: Ceiling division threads over-launch karta hai; check extras ko idle karta hai taaki woh out of bounds na likhein.