Interleaved — Phase 6

Hardware interleaved practice

printable — key stays hidden on paper

Parallel Architectures & GPU Computing

Instructions: Solve each problem. Each mixes a different subtopic — decide the correct model/formula before computing. Show numeric work where asked. Total: 100 marks.

1. A program spends 30% of its runtime in a strictly sequential section; the remaining 70% is perfectly parallelizable. Using Amdahl's Law, compute the maximum achievable speedup as the number of processors $p \to \infty$ , and the actual speedup at $p = 8$ . (10 marks)

2. A CUDA kernel launches a grid of $256$ blocks, each with $512$ threads. Given a warp size of $32$ , compute (a) the total number of threads, (b) the number of warps per block, and (c) the total number of warps in the grid. (9 marks)

3. Classify each of the following under Flynn's taxonomy (SISD / SIMD / MIMD): (a) a classic single-core scalar CPU running one instruction stream; (b) a GPU applying one instruction to a vector of data lanes; (c) a dual-socket server running independent threads on different data. Justify each in one line. (9 marks)

4. In a shared-memory multiprocessor, two threads on different cores each repeatedly update two different int counters (4 bytes each) that happen to reside in the same 64-byte cache line. Explain the performance problem, name it, and state one code-level fix. (10 marks)

5. A weak-scaling workload keeps per-processor work fixed. Using Gustafson's Law with a serial fraction $s = 0.1$ (parallel fraction $0.9$ ), compute the scaled speedup at $p = 16$ . (9 marks)

6. A warp of 32 threads executes:

if (threadIdx.x % 2 == 0) { A(); } else { B(); }

where A() and B() each take the same time $t$ . Assuming SIMT lock-step execution, what is the effective execution time relative to a divergence-free warp, and why? Name the phenomenon. (9 marks)

7. Contrast a multicore (e.g., 8 fat out-of-order cores) design with a manycore/GPU (thousands of simple cores) design in terms of the parallelism each targets — connect your answer to instruction-level vs thread-level parallelism. (10 marks)

8. In a directory-based cache coherence protocol on a 64-node system, explain why a directory scheme is used instead of snooping, and describe what the directory entry tracks for one memory block. (10 marks)

9. Shared memory in an SM is organized into 32 banks. Thread $i$ in a warp accesses sdata[i * 2] (4-byte words). Determine whether bank conflicts occur, the conflict degree, and how to fix it. (9 marks)

10. A global-memory load has 32 threads of a warp accessing consecutive 4-byte elements data[tid] starting at an aligned address. State whether this access is coalesced, how many memory transactions (128-byte) it needs, and contrast with a strided access data[tid * 16]. (5 marks)

Answer keyMark scheme & solutions

1. (Amdahl's Law — 6.1.3) Formula: $S(p) = \dfrac{1}{(1-f) + f/p}$ , with parallel fraction $f = 0.7$ , serial $= 0.3$ .

$p \to \infty$ : $S = \dfrac{1}{0.3} = 3.33$ .
$p = 8$ : $S = \dfrac{1}{0.3 + 0.7/8} = \dfrac{1}{0.3 + 0.0875} = \dfrac{1}{0.3875} = 2.58$ .

Why this method: Fixed problem size + fixed serial fraction ⇒ Amdahl (strong scaling), not Gustafson.

2. (Thread blocks and grids / warps — 6.2.6, 6.2.5)

(a) $256 \times 512 = 131072$ threads.
(b) $512 / 32 = 16$ warps per block.
(c) $16 \times 256 = 4096$ warps.

Why: Pure CUDA execution-model bookkeeping; warp = 32-thread scheduling unit.

3. (Flynn's taxonomy — 6.1.1)

(a) SISD — one instruction stream, one data stream.
(b) SIMD — one instruction broadcast across many data lanes (GPU/vector; technically SIMT, a SIMD variant).
(c) MIMD — multiple independent instruction streams on multiple data (multiprocessor/multithread).

Why: Classification by number of concurrent instruction vs data streams.

4. (False sharing — 6.1.10) Problem: the two counters share a cache line, so each write invalidates the line in the other core's cache (coherence ping-pong), causing serialization and wasted bandwidth despite logically independent data. Name: false sharing. Fix: pad/align each counter to its own cache line (e.g., alignas(64) or add padding), or give each thread a private local copy and reduce at the end.

Why: Two independent variables + one line ⇒ false sharing, not a real data race.

5. (Gustafson's Law — 6.1.3) Formula: $S(p) = p - s(p-1)$ with $s = 0.1$ . $S = 16 - 0.1(15) = 16 - 1.5 = 14.5$ .

Why: Per-processor work fixed (weak scaling) ⇒ Gustafson, which scales the problem size with $p$ — contrast Problem 1 which fixed the problem.

6. (Warp divergence — 6.2.11 / SIMT — 6.2.4) Because SIMT executes a warp in lock-step, the if and else branches cannot run simultaneously when threads diverge. The hardware serializes them: even-lane threads run A() while odd lanes are masked off, then odd lanes run B() while even are masked. Effective time $\approx 2t$ (both branches executed sequentially) vs $t$ for a divergence-free warp. Phenomenon: warp/branch divergence.

Why: Intra-warp data-dependent branching ⇒ divergence penalty (predication/masking), not a scheduling win.

7. (Multicore vs manycore + ILP/TLP — 6.1.4, 6.1.2, 6.2.1)

Multicore fat cores: few cores with deep pipelines, out-of-order execution, large caches, branch prediction — optimized to extract instruction-level parallelism (ILP) from a single/few threads and minimize latency.
Manycore/GPU: thousands of simple in-order cores with massive hardware multithreading — optimized for thread-level (and data-level) parallelism (TLP), hiding latency via throughput rather than reducing it.

Why: Distinguishes latency-optimized ILP machines from throughput-optimized TLP machines.

8. (Directory-based coherence — 6.1.6) Snooping broadcasts every coherence transaction on a shared bus; at 64 nodes the broadcast traffic and bus bandwidth don't scale. A directory keeps coherence info in a distributed structure so messages are sent point-to-point only to sharers. Each directory entry per block tracks: the state (e.g., Modified/Shared/Invalid) and a sharer list / presence bit-vector (which nodes hold a copy) plus the owner for a dirty block.

Why: Scalability of coherence ⇒ directory replaces broadcast snooping.

9. (Bank conflicts — 6.2.9) Thread $i$ accesses word index $2i$ . Bank $= (2i) \bmod 32$ . For $i$ and $i+16$ : $2i \bmod 32 = 2(i+16) \bmod 32$ , so threads $i$ and $i+16$ hit the same bank ⇒ 2-way bank conflict across 16 pairs. Fix: pad the array stride to an odd value (e.g., use stride 33 / access sdata[i*2 + i]-style padding) or restructure so the stride is coprime with 32.

Why: Even stride into a power-of-two bank count ⇒ conflicts; only strides odd w.r.t. 32 are conflict-free.

10. (Coalesced access — 6.2.8) data[tid]: 32 threads × 4 bytes = 128 contiguous aligned bytes ⇒ fully coalesced, served by 1 128-byte transaction. data[tid*16]: each thread touches a different 64-byte segment ⇒ scattered, requiring up to 32 transactions — massively wasted bandwidth.

Why: Contiguous aligned access = coalesced; strided access breaks coalescing.

[
  {"claim":"Amdahl speedup at p=8 with f=0.7 is 1/0.3875",
   "code":"S=1/(0.3+0.7/8); result = abs(S-2.5806451612903225)<1e-9"},
  {"claim":"Gustafson scaled speedup s=0.1 p=16 equals 14.5",
   "code":"p=16; s=0.1; S=p-s*(p-1); result = S==14.5"},
  {"claim":"Total warps in grid = 4096",
   "code":"blocks=256; tpb=512; warp=32; result = blocks*(tpb//warp)==4096"}
]