GPU Architecture

Question 1 (10 marks)

(a) A warp = a group of 32 threads scheduled and executed together in lockstep by an SM (SIMT). All threads share one instruction stream. When threads in the same warp take different branch paths, the hardware must serialize the paths — executing path A with B-threads masked off, then path B with A-threads masked off. Threads split within a warp cause divergence. (3: warp=32 threads/lockstep 1, single instruction stream 1, serialization mechanism 1)

(b) Within one warp (lanes 0–31): lanes 0–15 take path A (20 instr), lanes 16–31 take path B (12 instr). Because both halves are in the same warp, they serialize: $20 + 12 = 32 \text{ cycles.}$ (Only convergent code runs at full efficiency; divergence cost = sum of both paths.) (3: identifies both halves in same warp 1, serialization sum 1, =32 1)

(c) Make the branch granularity a multiple of the warp size (32) so every warp is uniform. Use:

if ((threadIdx.x / 32) % 2 == 0)  // even warps -> A, odd warps -> B

or equivalently if (threadIdx.x < 128). Now every warp is entirely in one path, so no warp diverges; the 256-thread block still splits 128/128 (16 warps → 8 each). (4: warp-aligned condition 2, correct 16/16-per-warp or 128/128 split 1, explanation no intra-warp split 1)

Question 2 (12 marks)

(a) Coalescing = combining the individual memory accesses of a warp's 32 threads into the minimum number of aligned memory transactions. A[threadIdx.x] accesses 32 consecutive floats = 32×4 = 128 bytes, contiguous and aligned. Transactions are 32 bytes each: $128 / 32 = 4 \text{ transactions.}$ Fully coalesced, 100% efficiency. (4: definition 1, 128 bytes 1, /32 1, =4 1)

(b) A[threadIdx.x*2] touches floats at indices 0,2,4,…,62 spanning 64 floats = 256 bytes, but only 128 bytes are used. Transactions needed = $256/32 = 8$ (every 32-byte segment is touched by the strided pattern). Useful data = 128 bytes of 256 fetched: $\text{efficiency} = 128/256 = 50\%.$ So 8 transactions, half the efficiency of (a). (4: 8 transactions 2, 50% efficiency 2)

(c) Bank = (word_index) mod 32. Thread t accesses word 2t. Bank = (2t) mod 32.

• t=0→0, t=1→2, …, t=15→30, t=16→0, t=17→2, … So threads 0 and 16 both hit bank 0; each bank is hit by exactly 2 threads. $\Rightarrow \textbf{2-way bank conflict.}$ Derivation: stride-2 access with 32 banks → conflict degree = gcd-driven; here 32/gcd(2,32)=16 distinct banks, 32 threads / 16 banks = 2 per bank. (4: bank formula 1, mapping/collision 2, 2-way 1)

Question 3 (10 marks)

(a) Compute blocks allowed per SM by each limit:

• Registers: per block = 48 × 256 = 12288 regs → 65536/12288 = 5.33 → 5 blocks.
• Shared mem: 64 KB / 4 KB = 16 blocks.
• Threads: 2048 / 256 = 8 blocks.
• Max blocks: 16 blocks.

Limiting = min(5, 16, 8, 16) = 5 blocks (register-limited).

Active threads = 5 × 256 = 1280. Active warps = 1280/32 = 40, but SM max = 32 warps. So warps are capped at 32… but blocks give 5×8 = 40 warps > 32. Recheck: 5 blocks × 8 warps = 40 > 32 → the 32-warp cap binds too. So actual = floor(32/8)=4 blocks by warp count.

Effective blocks = min(5 reg, 4 warp-cap, 8, 16) = 4 blocks. Active warps = 4×8 = 32. Occupancy = 32/32 = 100%… but active threads limited by regs to 5 blocks worth; since only 4 blocks fit due to warp cap, occupancy = 32 active warps / 32 max = 100%.

Wait — resolve cleanly: register limit allows 5 blocks (=40 warps) but the SM caps at 32 warps, so only 4 blocks actually run. Occupancy = 32/32 = 100%.

(6: register calc 2, shared/thread limits 1, warp-cap consideration 2, final occupancy 1)

(b) Latency hiding: when a warp stalls on a long-latency operation (e.g. ~400-cycle global load), the SM's warp scheduler switches to another ready warp, keeping execution units busy. Higher occupancy = more warps available to hide latency. But peak throughput needs only enough warps to cover the latency; beyond that point, extra occupancy gives no benefit — and high occupancy may force fewer registers/less shared memory, hurting per-thread performance. Hence 100% is not always optimal. (4: scheduler swaps stalled warps 2, "enough not maximal" 2)

Question 4 (12 marks)

(a) Kernel with grid-stride loop:

__global__ void vadd(const float* A, const float* B, float* C, int N) {
    int idx    = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = idx; i < N; i += stride)
        C[i] = A[i] + B[i];
}

(5: global qualifier 1, global index 1, stride 1, loop 1, add 1)

(b) Host code:

float *dA, *dB, *dC;
size_t bytes = N * sizeof(float);
cudaMalloc(&dA, bytes);
cudaMalloc(&dB, bytes);
cudaMalloc(&dC, bytes);
cudaMemcpy(dA, hA, bytes, cudaMemcpyHostToDevice);
cudaMemcpy(dB, hB, bytes, cudaMemcpyHostToDevice);
 
int block = 256;
int grid  = (N + block - 1) / block;
vadd<<<grid, block>>>(dA, dB, dC, N);
cudaDeviceSynchronize();
 
cudaMemcpy(hC, dC, bytes, cudaMemcpyDeviceToHost);
cudaFree(dA); cudaFree(dB); cudaFree(dC);

(5: malloc 1, H2D 1, grid calc + launch 2, D2H + free 1)

(c) A grid-stride loop lets a fixed, tuned grid size handle any N, decoupling launch config from data size. It improves reuse (better instruction/thread efficiency), tolerates N not divisible by block size, and lets each thread process multiple elements — improving memory-level parallelism and occupancy tuning. (2)

Question 5 (10 marks)

(a)

Order: registers (fastest) > shared > global. (3: correct ranking 1.5, latency magnitudes 1.5)

(b) Ridge-point arithmetic intensity: $AI_{ridge} = \frac{\text{peak FLOP/s}}{\text{peak bytes/s}} = \frac{15\times10^{12}}{900\times10^{9}} = 16.67 \text{ FLOP/byte.}$ Kernel AI = 0.25 FLOP/byte ≪ 16.67 → memory-bound. (4: ridge formula 1, =16.67 1, compare 1, memory-bound 1)

(c) Any two: (1) ensure coalesced global accesses; (2) use shared memory to reuse data (tiling); (3) vectorized loads (float4); (4) increase occupancy for latency hiding; (5) minimize redundant global traffic / fuse kernels. (3: 1.5 each)

Question 6 (6 marks)

(a) Tensor cores accelerate small matrix-multiply-accumulate: $D = A\times B + C$ on tiles (e.g. 4×4 or 16×16) in a single instruction. Faster than CUDA-core FMAs because one instruction performs many MACs in parallel with dedicated hardware and lower per-op overhead, exploiting mixed-precision (FP16/BF16 input, FP32 accumulate) for higher throughput. (3: MMA operation 1.5, parallel/dedicated HW reason 1.5)

(b) AMD equivalent = ROCm/HIP. Portable standard running on both = OpenCL (also SYCL). (3: ROCm/HIP 1.5, OpenCL 1.5)

[
  {"claim":"Q1b divergence cycles = 20+12 = 32","code":"result = (20+12)==32"},
  {"claim":"Q2a coalesced 128 bytes / 32-byte transactions = 4","code":"result = (32*4)//32==4"},
  {"claim":"Q2b strided efficiency 128/256 = 50%","code":"result = sympy.Rational(128,256)==sympy.Rational(1,2)"},
  {"claim":"Q2c 32 banks stride-2 => 2-way conflict","code":"banks=set([(2*t)%32 for t in range(32)]); result = (32//len(banks))==2"},
  {"claim":"Q3 register limit floor(65536/(48*256)) = 5 blocks","code":"result = 65536//(48*256)==5"},
  {"claim":"Q3 warp cap: 5 blocks*8 warps=40>32 so 4 blocks fit","code":"result = (32//8)==4 and 5*8>32"},
  {"claim":"Q5b roofline ridge AI = 15e12/900e9 approx 16.67","code":"result = abs(15e12/900e9 - 16.666666666)<1e-3"},
  {"claim":"Q4b grid for N and block 256 uses ceil division","code":"N=1000; block=256; result = (N+block-1)//block==4"}
]