CUDA programming model basics: Exercises — Hinglish deep dive

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

Shuru karne se pehle, ek formula jis par poora page tika hai usse pakka kar lete hain, kyunki isme har symbol ko samajhna zaroori hai.

Figure dekho: crate $2$ , seat $3$ , $4$ lockers per crate ke saath, locker $2 \times 4 + 3 = 11$ par land karta hai.

L1 — Recognition

Exercise 1.1 (L1)

Har term ko uski jagah se match karo. Host, Device, Kernel, Warp ke liye ek phrase mein batao yeh kya hai.

Recall Solution

Host ::: CPU aur uski system RAM — jahan main() run karta hai.
Device ::: GPU aur uski VRAM — jahan kernels run karte hain.
Kernel ::: ek __global__ function jo host se launch hota hai lekin device par bahut saare threads ke saath ek saath execute hota hai.
Warp ::: $32$ threads ka ek hardware bundle jo instructions ko lockstep (SIMT) mein execute karta hai. Dekho Thread Warps and SIMT.

Exercise 1.2 (L1)

vectorAdd<<<numBlocks, threadsPerBlock>>>(...) mein, donon triple-angle numbers mein se kaun sa grid size hai aur kaun sa block size?

Recall Solution

Pehla number, numBlocks ::: grid size — kitne blocks hain.
Doosra number, threadsPerBlock ::: block size — har block mein kitne threads hain. Toh logical order hai Grid → Blocks → Threads, bilkul waisi hi order jisme yeh appear karte hain.

Exercise 1.3 (L1)

Inhe fastest→slowest order mein lagao: Global Memory, Registers, Shared Memory, Host Memory.

Recall Solution

Registers ( $\approx 1$ cycle) → Shared Memory ( $\approx 5$ cycles) → Global Memory ( $\approx 200 - 400$ cycles) → Host Memory ( $\approx 500 - 1000$ cycles). Compute cores ke jitna paas = utne kam wires cross karne padte hain = utna fast. Dekho Memory Hierarchy.

L2 — Application

Exercise 2.1 (L2)

Ek thread ka blockIdx.x = 3, blockDim.x = 256, threadIdx.x = 17 hai. Uska globalIdx compute karo.

Recall Solution

$globalIdx = 3 \times 256 + 17 = 768 + 17 = 785.$ Humne kya kiya: crates $0, 1, 2$ mein $768$ threads count kiye, phir apni seat $17$ add ki.

Exercise 2.2 (L2)

Tumhare paas $N = 1, 000, 000$ elements hain aur tum $threadsPerBlock = 256$ choose karte ho. Kitne blocks launch karoge?

Recall Solution

Ceiling division use karo taaki koi element cover hone se na reh jaye: $numBlocks = ⌈ \frac{N}{256} ⌉ = \frac{1000000 + 256 - 1}{256} = \frac{1000255}{256} = 3907.$ Check: $3907 \times 256 = 1, 000, 192 \geq 1, 000, 000$ . Theek hai — kuch spare threads hain, jinhe boundary check handle karta hai.

Exercise 2.3 (L2)

width = 640 ki ek $2$ D image ke liye, ek thread row = 5, col = 12 compute karta hai. Row-major array mein uska flattened globalIdx kya hoga?

Recall Solution

Row-major ka matlab hai: agli row shuru karne se pehle har completed row ko poora cross karo. $globalIdx = row \times width + col = 5 \times 640 + 12 = 3200 + 12 = 3212.$

L3 — Analysis

Exercise 3.1 (L3)

$N = 1000$ aur $threadsPerBlock = 256$ ke saath, exactly kitne threads launch hote hain, aur if (idx < N) ki wajah se kitne threads idle hain (jo koi kaam nahi karte)?

Recall Solution

$numBlocks = ⌈ 1000/256 ⌉ = 4$ . Total threads launched $= 4 \times 256 = 1024$ . Working threads = jinke $idx < 1000$ hai, yaani unme se $1000$ . Idle threads $= 1024 - 1000 = 24$ (indices $1000$ se $1023$ tak). Yeh if test ko false pass karte hain aur simply exit kar jaate hain.

Exercise 3.2 (L3)

Figure dekho. Tail crate sirf partly use ho rahi hai. Explain karo kyun vectorAdd kernel se if (idx < N) hatana dangerous hai, aur $N = 1000$ ke liye kaun sa thread sabse pehle misbehave karta hai.

Recall Solution

Threads poore blocks mein aate hain; tum "exactly $1000$ " launch nahi kar sakte. Last block index $1023$ tak overshoot karta hai. Guard ke bina, thread $idx = 1000$ A[1000]/B[1000] read karega aur C[1000] write karega — array ke end se ek element aage (valid indices $999$ par ruk jaate hain). Yeh out-of-bounds access hai: crash ya corrupted neighbouring memory. Pehla offender $idx = 1000$ hai.

Exercise 3.3 (L3)

Ek math op $\approx 1$ cycle cost karta hai; ek global memory read $\approx 300$ cycles cost karta hai. Naïve vectorAdd mein, har output ko $2$ global reads + $1$ global write + $1$ add chahiye. Actual arithmetic par roughly kitna fraction of time spend hota hai?

Recall Solution

Memory cycles $\approx 3 \times 300 = 900$ . Arithmetic $\approx 1$ cycle. $arithmetic fraction = \frac{1}{900 + 1} = \frac{1}{901} \approx 0.00111 = 0.111%.$ Interpretation: vectorAdd bilkul memory-bound hai — GPU apna $99.9%$ time data move karne mein spend karta hai. Isliye Memory Hierarchy aur coalescing raw FLOPs se zyada matter karte hain yahan.

L4 — Synthesis

Exercise 4.1 (L4)

Ek kernel scaleAdd likho jo $C [i] = a \cdot A [i] + B [i]$ compute kare (ek "SAXPY") proper index aur boundary check ke saath. $N = 500000$ , $threadsPerBlock = 256$ ke liye launch config batao.

Recall Solution

__global__ void scaleAdd(float a, float *A, float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = a * A[idx] + B[idx];
    }
}

Launch: $numBlocks = ⌈ \frac{500000}{256} ⌉ = \frac{500000 + 255}{256} = \frac{500255}{256} = 1954.$ Check: $1954 \times 256 = 500224 \geq 500000$ . Call: scaleAdd<<<1954,256>>>(a,d_A,d_B,d_C,N);

Exercise 4.2 (L4)

Kisi bhi kernel ko device data par run karne ke liye zaroori chhe host-side steps (order mein) list karo, aur input aur output copies ke liye correct cudaMemcpy direction enum do.

Recall Solution

Host memory allocate karo (malloc).
Device memory allocate karo (cudaMalloc).
Inputs Host→Device copy karo — enum cudaMemcpyHostToDevice.
Kernel <<<numBlocks, threadsPerBlock>>> launch karo.
Results Device→Host copy karo — enum cudaMemcpyDeviceToHost.
Device (cudaFree) aur host (free) memory free karo.

Exercise 4.3 (L4)

$640 \times 480$ image (width=640, height=480) par $2$ D kernel ke liye block dims $16 \times 16$ ke saath, grid dimensions (gridX, gridY) compute karo.

Recall Solution

$gridX = ⌈ \frac{640}{16} ⌉ = 40, gridY = ⌈ \frac{480}{16} ⌉ = 30.$ Total blocks $= 40 \times 30 = 1200$ ; total threads $= 1200 \times 256 = 307200$ , jo $640 \times 480 = 307200$ pixels se exactly match karta hai (yahan koi idle threads nahi kyunki dono dims evenly divide hote hain).

L5 — Mastery

Exercise 5.1 (L5)

Ek kernel grid $= 3907$ blocks, block $= 256$ threads ke saath launch hota hai. Jo thread array index $200000$ process karne wala hai — uske blockIdx.x aur threadIdx.x kya hain?

Recall Solution

Index formula ko invert karo. Block size se divide karo: $blockIdx.x = ⌊ \frac{200000}{256} ⌋ = 781, threadIdx.x = 200000 mod 256 = 200000 - 781 \times 256.$ $781 \times 256 = 199936$ , toh $threadIdx.x = 200000 - 199936 = 64$ . Forward verify karo: $781 \times 256 + 64 = 199936 + 64 = 200000$ . ✓

Exercise 5.2 (L5)

Maano global memory se read karna $300$ cycles cost karta hai lekin shared memory se $5$ cycles. Ek tiling kernel har element ko global memory se ek baar shared memory mein load karta hai, phir usse shared memory se $8$ baar reuse karta hai. Naïve kernel ke against total memory cycles per element compare karo jo global memory $8$ baar read karta hai.

Recall Solution

Naïve: $8$ global reads $= 8 \times 300 = 2400$ cycles.
Tiled: $1$ global read $+ 8$ shared reads $= 300 + 8 \times 5 = 300 + 40 = 340$ cycles.
Speedup $= 2400/340 \approx 7.06 \times$ . Isliye exactly shared memory exist karti hai: expensive off-chip trip ek baar pay karo, phir saste mein reuse karo. Dekho Memory Hierarchy aur Streaming Multiprocessors.

Exercise 5.3 (L5)

Ek GPU ke paas $80$ SMs hain aur maximum $8$ resident blocks per SM hold kar sakta hai. Tum $3907$ blocks ka ek grid launch karte ho. Kya saare blocks ek saath resident hain? Agar nahi, toh GPU ko kitne "waves" of blocks ki zaroorat hai?

Recall Solution

Max simultaneously resident blocks $= 80 \times 8 = 640$ . Kyunki $3907 > 640$ hai, saare ek saath fit nahi hote. Number of waves: $⌈ \frac{3907}{640} ⌉ = \frac{3907 + 639}{640} = \frac{4546}{640} = 7.10 \dots \Rightarrow 7 waves (rounded up) .$ Actually $⌈ 3907/640 ⌉ = ⌈ 6.10 ⌉ = 7$ . GPU blocks ko $7$ waves mein process karta hai; scheduler finished blocks retire karta hai aur waiting wale swap in karta hai. Yahi cheez ek kernel ko code change kiye bina ek small GPU se huge GPU tak scale karne deti hai.

Recall Quick self-test recap

N=1000, tpb=256 ke liye ceiling block count ::: $⌈ 1000/256 ⌉ = 4$ blocks, $24$ idle threads. Block 3, blockDim 256, thread 17 ke liye global index ::: $3 \times 256 + 17 = 785$ . if (idx < N) kyun ::: last block array ko overshoot karta hai; guard out-of-bounds access rokta hai. Fastest CUDA memory ::: registers ( $\approx 1$ cycle), on-chip, compute unit mein.