• SIMT model: threads, warps/wavefronts
• SM/CU microarchitecture
• Memory hierarchy and coalescing
• Warp scheduling and latency hiding
• Control-flow divergence
• Occupancy and resource limits
• Kernels, grids/blocks, synchronization
• Performance tuning patterns
• Exercises

• GPUs execute groups of threads in lockstep (NVIDIA: 32-thread warps; AMD: 32/64 wavefronts).
• Same instruction across lanes with per-lane masks; hardware manages active masks.
• Latency hiding via switching among many resident warps.

Warp: 32 lanes → issue 1 inst/cycle (idealized)
Mask: 111111... → lanes disabled on divergence

• Streaming Multiprocessor (SM, NVIDIA) / Compute Unit (CU, AMD) houses vector ALUs, load/store units, SFUs, and schedulers.
• Instruction issue from multiple warp schedulers; dual-issue on certain mixes.
• Register file and shared memory are partitioned per SM; limits occupancy.

• Global (device) memory with high bandwidth (GDDR/HBM), higher latency.
• L2 and per-SM L1/texture caches; software-managed shared memory (scratchpad).
• Coalescing: adjacent threads should access adjacent addresses for efficient bursts.

// Coalesced pattern (CUDA-style pseudo)
int i = blockIdx.x * blockDim.x + threadIdx.x;
a[i] = b[i] + c[i];

• Fine-grained thread scheduling selects a ready warp each cycle.
• Long-latency ops (global loads, SFU) are hidden by switching to other ready warps.
• Issue policies vary (round-robin, GTO); backpressure from memory can throttle.

• Divergent branches split warp execution into serialized paths with masks.
• Structure kernels to minimize divergence within warps (e.g., partition data).
• Use predication or warp-level primitives to reduce control overheads.

• Occupancy = resident warps per SM limited by registers/thread and shared memory/block.
• Higher occupancy aids latency hiding but can increase pressure and reduce per-thread resources.
• Balance block size, registers, and shared memory to achieve good occupancy.

• Hierarchy: grid → blocks → threads; blocks scheduled to SMs.
• Intra-block sync via barriers; inter-block sync requires kernel boundaries or cooperative groups.
• Warp-level intrinsics (shuffles, ballots) enable efficient collectives.

• Tile into shared memory to improve locality; avoid bank conflicts.
• Ensure memory coalescing; use vectorized loads/stores where appropriate.
• Use occupancy calculators; profile with Nsight/rocprof to find bottlenecks.

1. Implement a tiled matrix multiply using shared memory; measure speedup vs naive.
2. Optimize a memory-bound kernel by improving coalescing and removing divergence.
3. Experiment with block sizes to study occupancy vs performance.