GPU architecture — SIMT, warps, CUDA model: Exercises

This page is a self-test ladder for GPU architecture — SIMT, warps, CUDA model. Work each problem before opening its solution. Every symbol used here was built in the parent note; if you get stuck on a term, that note defines it.

Prerequisite side-links you may want open: SIMD and Vector Instructions (AVX, SSE), Memory Coalescing and Bandwidth, CUDA Shared Memory and Tiling, Amdahl's and Gustafson's Laws, Matrix Multiplication on GPUs, CPU Cache Hierarchy and Latency.

Level 1 — Recognition

L1.1 — Warp size

Problem. On an NVIDIA GPU, how many threads run in lockstep inside one warp, and what does "lockstep" mean in one sentence?

Recall Solution

A warp is 32 threads. "Lockstep" means: one instruction is fetched and decoded once, then all 32 threads execute that same instruction in the same cycle (each on its own data / registers). They advance together, one instruction at a time.

L1.2 — Name the level of the hierarchy

Problem. Match each name to what it groups: thread, warp, block, grid.

Recall Solution

thread — one execution of the kernel, has a unique index.
warp — a hardware bundle of 32 consecutive threads (the scheduling unit).
block (CTA) — a group of threads that share on-chip shared memory and can __syncthreads(); runs entirely on one SM.
grid — all blocks of one kernel launch.

L1.3 — `blockDim` vs `gridDim`

Problem. In CUDA, one of blockDim.x and gridDim.x is "threads per block" and the other is "number of blocks." Which is which?

Recall Solution

blockDim.x = threads per block.
gridDim.x = number of blocks in the grid. Memory hook: Dim = the size of the thing named before it. blockDim = size of a block (in threads).

Level 2 — Application

L2.1 — Global thread index

Problem. A kernel launches with blockDim.x = 256. The thread with blockIdx.x = 5 and threadIdx.x = 17 runs. Which global element index $i$ does it own?

Recall Solution

Use the derived formula $i = blockIdx.x \cdot blockDim.x + threadIdx.x = 5 \cdot 256 + 17 = 1280 + 17 = 1297.$ Read it as "skip the 5 earlier blocks (each 256 wide), then step 17 into my own block."

L2.2 — Number of blocks (ceiling division)

Problem. You have $N = 1000$ elements and choose block size $B = 256$ . How many blocks must you launch so no element is dropped?

Recall Solution

$numBlocks = ⌈ \frac{N}{B} ⌉ = \frac{N + B - 1}{B} = \frac{1000 + 255}{256} = \frac{1255}{256} = 4 (integer division) .$ Check: 4 blocks × 256 = 1024 threads ≥ 1000. Three blocks would give only 768 < 1000, dropping elements. So 4 blocks, and you add the guard if (i >= N) return; so the 24 extra threads (indices 1000–1023) exit before touching memory.

L2.3 — Wasted lanes

Problem. A block has 40 threads. How many warps does it occupy, and how many lanes sit idle?

Recall Solution

Warps per block $= ⌈ 40/32 ⌉ = 2$ . Those 2 warps supply $2 \times 32 = 64$ lane-slots, but only 40 threads exist, so $64 - 40 = 24$ lanes are idle. The second warp holds threads 32–39 (8 active) and 24 masked-off lanes.

Level 3 — Analysis

L3.1 — Coalescing efficiency for a strided read

Problem. A warp of 32 threads each reads one 4-byte float. Memory is served in 128-byte transactions. Compute the coalescing efficiency for (a) stride 1 (a[i]) and (b) stride 16 (a[i*16]).

Recall Solution

Efficiency $= \frac{bytes requested}{bytes actually transferred}$ . Bytes requested $= 32 \times 4 = 128$ B in both cases. (a) stride 1: the 32 addresses span $32 \times 4 = 128$ B — exactly one 128-byte line. Transferred $= 128$ B. $eff = \frac{128}{128} = 1.$ (b) stride 16: consecutive threads are $16 \times 4 = 64$ B apart. Thread $t$ touches byte $64 t$ ; the 32 touched bytes span $64 \times 31 = 1984$ B, landing in $⌈ 1984/128 ⌉ + alignment$ ... more simply, each thread lands in a different 128-byte line once the stride ( $64$ B) is at least half a line and offsets don't share lines: here two consecutive threads (64 B apart) fall into the same 128 B line only in pairs, so lines touched $= 16$ , transferred $= 16 \times 128 = 2048$ B. $eff = \frac{128}{2048} = \frac{1}{16} .$ The picture (see figure) makes it obvious: stride-1 packs into one line, stride-16 scatters across many.

Figure — GPU architecture — SIMT, warps, CUDA model

L3.2 — Divergence cost

Problem. A warp hits if (threadIdx.x % 2 == 0) heavy_A(); else heavy_B();. Suppose heavy_A costs 30 cycles and heavy_B costs 50 cycles. What does this warp cost, and what would a warp-uniform branch cost?

Recall Solution

Even lanes want heavy_A, odd lanes want heavy_B — the same warp contains both, so both paths run serially with masking: $T_{warp} = T_{A} + T_{B} = 30 + 50 = 80 cycles .$ If instead all 32 lanes of the warp agreed (e.g. branch on threadIdx.x / 32), only one path runs: $T_{uniform} = max (T_{A}, T_{B}) = max (30, 50) = 50 cycles .$ Divergence overhead here $= 80 - 50 = 30$ cycles, a $80/50 = 1.6 \times$ slowdown.

L3.3 — Latency hiding: how many warps?

Problem. A global load takes 400 cycles; each warp does 5 cycles of independent arithmetic before its next load. How many warps must an SM keep resident to fully hide the 400-cycle latency (assume the scheduler issues one warp's instruction per cycle)?

Recall Solution

While one warp is stalled 400 cycles, the scheduler needs other warps to fill those cycles. Each warp offers 5 cycles of work between stalls. To cover 400 cycles you need $warps \geq \frac{400}{5} = 80.$ So $\geq 80$ resident warps keep every cycle busy — the 400-cycle latency is fully hidden. Fewer than 80 leaves bubbles where the ALUs idle.

Level 4 — Synthesis

L4.1 — Choose a block size

Problem. An SM supports at most 2048 resident threads and at most 32 resident blocks. You want maximum occupancy (most resident threads). Compare block sizes $B = 64$ and $B = 256$ . Which achieves full 2048-thread occupancy, and why can the smaller one fail?

Recall Solution

Resident threads $= min (32 blocks \times B, 2048)$ .

$B = 64$ : block-limit gives $32 \times 64 = 2048$ ... equal to the thread cap, so 2048 threads — full. But it uses all 32 block slots, leaving no headroom, and any per-block resource (shared memory) is multiplied by 32 blocks.
$B = 256$ : $2048/256 = 8$ blocks fit within both the 32-block cap and the 2048-thread cap → 2048 threads with only 8 blocks.

Both reach 2048 threads here, but $B = 64$ is fragile: if the block cap were 16 instead of 32, then $B = 64$ tops out at $16 \times 64 = 1024$ threads (half occupancy), while $B = 256$ still gives $min (16 \times 256, 2048) = 2048$ . Lesson: small blocks can hit the block-count ceiling before the thread ceiling.

L4.2 — Coalesced matrix access pattern

Problem. A row-major matrix $A$ has element $A [r] [c]$ at linear address $r \cdot W + c$ (width $W$ ). A warp of 32 threads processes 32 elements. To read coalesced, should the 32 threads of a warp vary the column $c$ (fixed row) or the row $r$ (fixed column)? Compute efficiency for each ( $W = 1024$ , 4-byte floats, 128-byte lines).

Recall Solution

Consecutive threads should hit consecutive addresses. Address $= (r \cdot W + c) \times 4$ bytes.

Vary $c$ (thread $t$ reads $A [r] [t]$ ): addresses are $4 \cdot (r W + t)$ — spacing 4 B, consecutive. The 32 span $128$ B → one line. $eff = 128/128 = 1$ . ✅
Vary $r$ (thread $t$ reads $A [t] [c]$ ): addresses are $4 \cdot (t W + c)$ — spacing $4 W = 4096$ B. Each thread lands in a different 128-byte line → 32 lines transferred $= 32 \times 128 = 4096$ B. $eff = 128/4096 = 1/32$ . ❌

So let the warp vary the fast-changing (column) index — matches the parent's stride-1 rule. Column-major access needs shared-memory tiling to fix.

Level 5 — Mastery

L5.1 — The bounds-guard direction

Problem. For $N = 1000$ , $B = 256$ you launch 4 blocks = 1024 threads. Someone writes the guard as if (i < N) return;. Exactly which threads survive to run the body, and why is this catastrophically wrong?

Recall Solution

return is an early exit: whoever satisfies the condition leaves. With if (i < N) return;, threads with $i < 1000$ (the valid ones, indices 0–999) all exit — and only $i \geq 1000$ (indices 1000–1023, the out-of-range ones) survive to run the body. That is the exact inversion of what you want: 1000 good threads do nothing, 24 bad threads read/write past the array → wrong results and out-of-bounds memory access. Correct guard: if (i >= N) return; — exit the invalid threads, keep the valid ones.

L5.2 — Amdahl ceiling on a partly serial kernel

Problem. A GPU program spends a fraction $p = 0.95$ of its work in a perfectly parallel kernel and $1 - p = 0.05$ in a serial host step. With effectively "infinite" parallel speedup on the parallel part, what is the maximum overall speedup (Amdahl's law)?

Recall Solution

Amdahl's law with speedup $s \to \infty$ on the parallel fraction: $S_{m a x} = \frac{1}{( 1 - p ) + \frac{p}{s}} s \to \infty \frac{1}{1 - p} = \frac{1}{0.05} = 20.$ No matter how many warps you throw at it, the 5% serial part caps the whole program at 20×. This is why GPU work is designed to shrink the serial fraction (fuse kernels, keep data on the device), not just add threads.

L5.3 — Divergence inside a reduction (edge case)

Problem. A tree reduction halves the number of active threads each step: if (threadIdx.x < stride) a[t] += a[t+stride]; with stride = 256, 128, ..., 1. For a block of 256 threads (8 warps of 32), at which strides does intra-warp divergence first appear, and why do the early steps NOT diverge?

Recall Solution

A warp is threads $32 k \dots 32 k + 31$ . The branch is threadIdx.x < stride.

stride ≥ 32 (256, 128, 64, 32): each warp is entirely below the threshold or entirely at/above it — e.g. at stride 128, warps 0–3 (threads 0–127) are all active, warps 4–7 all inactive. No warp is split → no divergence.
stride = 16 (and 8, 4, 2, 1): now the threshold cuts through warp 0 (threads 0–15 active, 16–31 inactive) → that warp diverges: both the active and idle "paths" serialize (idle lanes masked). Divergence first appears at stride 16.

So the classic optimisation is to keep strides $\geq 32$ warp-uniform, then handle the final $< 32$ portion with warp-level primitives (shuffle) to avoid the divergent tail.

Recall One-line self-quiz (cover the answers)

Warp size ::: 32 threads in lockstep. Global index formula ::: $i = blockIdx \cdot blockDim + threadIdx$ . numBlocks for N=1000, B=256 ::: $⌈ 1000/256 ⌉ = 4$ . Correct bounds guard ::: if (i >= N) return; (exit the invalid threads). Divergent-warp cost of two paths ::: $T_{A} + T_{B}$ , vs $max (T_{A}, T_{B})$ if uniform. Amdahl ceiling at p=0.95 ::: $1/ (1 - p) = 20 \times$ .