CUDA programming model basics: Exercises — Hardware deep dive

Before we start, let us lock down the one formula the whole page leans on, because every symbol in it must be earned.

Look at the figure: crate $2$ , seat $3$ , with $4$ lockers per crate, lands on locker $2 \times 4 + 3 = 11$ .

L1 — Recognition

Exercise 1.1 (L1)

Match each term to its home. For Host, Device, Kernel, Warp, say in one phrase what it is.

Recall Solution

Host ::: the CPU and its system RAM — where main() runs.
Device ::: the GPU and its VRAM — where kernels run.
Kernel ::: a __global__ function launched from the host but executed on the device by many threads at once.
Warp ::: a hardware bundle of $32$ threads that step through instructions in lockstep (SIMT). See Thread Warps and SIMT.

Exercise 1.2 (L1)

In vectorAdd<<<numBlocks, threadsPerBlock>>>(...), which of the two triple-angle numbers is the grid size and which is the block size?

Recall Solution

First number, numBlocks ::: the grid size — how many blocks.
Second number, threadsPerBlock ::: the block size — how many threads per block. So the logical order is Grid → Blocks → Threads, exactly the order they appear.

Exercise 1.3 (L1)

Order these memories fastest→slowest: 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). Closer to the compute cores = fewer wires to cross = faster. See Memory Hierarchy.

L2 — Application

Exercise 2.1 (L2)

A thread has blockIdx.x = 3, blockDim.x = 256, threadIdx.x = 17. Compute its globalIdx.

Recall Solution

$globalIdx = 3 \times 256 + 17 = 768 + 17 = 785.$ What we did: counted the $768$ threads in crates $0, 1, 2$ , then added our seat $17$ .

Exercise 2.2 (L2)

You have $N = 1, 000, 000$ elements and choose $threadsPerBlock = 256$ . How many blocks do you launch?

Recall Solution

Use ceiling division so no element is left uncovered: $numBlocks = ⌈ \frac{N}{256} ⌉ = \frac{1000000 + 256 - 1}{256} = \frac{1000255}{256} = 3907.$ Check: $3907 \times 256 = 1, 000, 192 \geq 1, 000, 000$ . Good — a few spare threads, which the boundary check handles.

Exercise 2.3 (L2)

For a $2$ D image of width = 640, a thread computes row = 5, col = 12. What is its flattened globalIdx in a row-major array?

Recall Solution

Row-major means: walk all the way across each finished row before starting the next. $globalIdx = row \times width + col = 5 \times 640 + 12 = 3200 + 12 = 3212.$

L3 — Analysis

Exercise 3.1 (L3)

With $N = 1000$ and $threadsPerBlock = 256$ , exactly how many threads are launched, and how many are idle (do no work) thanks to if (idx < N)?

Recall Solution

$numBlocks = ⌈ 1000/256 ⌉ = 4$ . Total threads launched $= 4 \times 256 = 1024$ . Working threads = the ones with $idx < 1000$ , i.e. $1000$ of them. Idle threads $= 1024 - 1000 = 24$ (indices $1000$ through $1023$ ). They pass the if test as false and simply exit.

Exercise 3.2 (L3)

Look at the figure. The tail crate is only partly used. Explain why removing if (idx < N) from the vectorAdd kernel is dangerous, and which thread first misbehaves for $N = 1000$ .

Recall Solution

Threads come in whole blocks; you cannot launch "exactly $1000$ ." The last block overshoots to index $1023$ . Without the guard, thread $idx = 1000$ would read A[1000]/B[1000] and write C[1000] — one element past the array end (valid indices stop at $999$ ). That is an out-of-bounds access: crash or corrupted neighbouring memory. The first offender is $idx = 1000$ .

Exercise 3.3 (L3)

A math op costs $\approx 1$ cycle; a global memory read costs $\approx 300$ cycles. In naïve vectorAdd, each output needs $2$ global reads + $1$ global write + $1$ add. Roughly what fraction of the time is spent on the actual arithmetic?

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 is utterly memory-bound — the GPU spends $99.9%$ of its time moving data. This is why Memory Hierarchy and coalescing matter more than raw FLOPs here.

L4 — Synthesis

Exercise 4.1 (L4)

Write a kernel scaleAdd computing $C [i] = a \cdot A [i] + B [i]$ (a "SAXPY") with a proper index and boundary check. State the launch config for $N = 500000$ , $threadsPerBlock = 256$ .

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)

List the six host-side steps (in order) needed to run any kernel on device data, and give the correct cudaMemcpy direction enum for input and output copies.

Recall Solution

Allocate host memory (malloc).
Allocate device memory (cudaMalloc).
Copy inputs Host→Device — enum cudaMemcpyHostToDevice.
Launch kernel <<<numBlocks, threadsPerBlock>>>.
Copy results Device→Host — enum cudaMemcpyDeviceToHost.
Free device (cudaFree) and host (free) memory.

Exercise 4.3 (L4)

For a $2$ D kernel on a $640 \times 480$ image (width=640, height=480) with block dims $16 \times 16$ , compute the grid dimensions (gridX, gridY).

Recall Solution

$gridX = ⌈ \frac{640}{16} ⌉ = 40, gridY = ⌈ \frac{480}{16} ⌉ = 30.$ Total blocks $= 40 \times 30 = 1200$ ; total threads $= 1200 \times 256 = 307200$ , which matches $640 \times 480 = 307200$ pixels exactly (no idle threads here since both dims divide evenly).

L5 — Mastery

Exercise 5.1 (L5)

A kernel launches with grid $= 3907$ blocks, block $= 256$ threads. The thread that should process array index $200000$ — what are its blockIdx.x and threadIdx.x?

Recall Solution

Invert the index formula. Divide by block size: $blockIdx.x = ⌊ \frac{200000}{256} ⌋ = 781, threadIdx.x = 200000 mod 256 = 200000 - 781 \times 256.$ $781 \times 256 = 199936$ , so $threadIdx.x = 200000 - 199936 = 64$ . Verify forward: $781 \times 256 + 64 = 199936 + 64 = 200000$ . ✓

Exercise 5.2 (L5)

Suppose reading from global memory costs $300$ cycles but from shared memory costs $5$ cycles. A tiling kernel loads each element from global memory once into shared memory, then reuses it $8$ times from shared memory. Compare total memory cycles per element against a naïve kernel that reads global memory all $8$ times.

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$ . This is exactly why shared memory exists: pay the expensive off-chip trip once, then reuse cheaply. See Memory Hierarchy and Streaming Multiprocessors.

Exercise 5.3 (L5)

A GPU has $80$ SMs and can hold at most $8$ resident blocks per SM. You launch a grid of $3907$ blocks. Are all blocks resident at once? If not, how many "waves" of blocks does the GPU need?

Recall Solution

Max simultaneously resident blocks $= 80 \times 8 = 640$ . Since $3907 > 640$ , not all fit at once. 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$ . The GPU processes blocks in $7$ waves; the scheduler retires finished blocks and swaps in waiting ones. This is what lets one kernel scale from a small GPU to a huge one without changing code.

Recall Quick self-test recap

Ceiling block count for N=1000, tpb=256 ::: $⌈ 1000/256 ⌉ = 4$ blocks, $24$ idle threads. Global index for block 3, blockDim 256, thread 17 ::: $3 \times 256 + 17 = 785$ . Why if (idx < N) ::: last block overshoots array; guard stops out-of-bounds access. Fastest CUDA memory ::: registers ( $\approx 1$ cycle), on-chip, in the compute unit.