GPU Architecture and Parallel Computing Concepts
Execution model
A GPU executes a kernel as a grid of thread blocks; each block runs on one SM and contains warps of 32 threads. Threads in a warp execute in lockstep on SIMT hardware, so divergent branches within a warp serialize. Candidates should be able to trace a thread's identity via blockIdx, threadIdx, and grid dimensions, and compute a global index correctly.
The memory hierarchy
Registers are fastest and per-thread; shared memory is per-block and programmer-managed; L1/L2 caches sit between; global memory (HBM) is large but high-latency. A frequent exam trap is treating shared memory as automatically beneficial: it helps only when it removes redundant global accesses and avoids bank conflicts. Coalesced global access, where consecutive threads read consecutive addresses, is the single most important performance habit.
Occupancy and latency hiding
Occupancy is the ratio of active warps to the maximum supported per SM. It is limited by registers, shared memory, and block size. High occupancy is not a goal by itself; enough parallelism to hide memory latency is. Traps include assuming maximum occupancy always wins and ignoring register pressure from large per-thread arrays.
Synchronization
__syncthreads() synchronizes threads within a block only, never across blocks. Cross-block coordination requires separate kernel launches, atomics, or cooperative groups. Warp-level primitives such as shuffle avoid shared memory round-trips.
Worked example
A kernel where each thread sums a strided column of a matrix is slow because accesses are uncoalesced. Fix: have consecutive threads read consecutive row elements and accumulate into shared memory tiles, then reduce. Rule of thumb: if adjacent threads touch addresses more than a cache line apart, restructure the access pattern first.