The GPU Execution Model
Assumed knowledge in every GPU interview. Draw the thread → warp → block → grid hierarchy, map it to SM hardware, and know why blocks have no ordering guarantee — that's the foundation everything else builds on.
Can you map the software hierarchy to hardware, know a block runs entirely on one SM, and explain why there's no ordering between blocks (it's what lets one binary scale across any GPU size)?
SIMT: Scalar Code, Lockstep Warps
You write scalar per-thread code; hardware groups threads into warps of 32 that issue in lockstep. The hierarchy: thread → warp (32) → block → grid. The scheduler assigns whole blocks to SMs; a block never spans two SMs.
So the runtime can schedule blocks in any order across any number of SMs — the same binary scales from a laptop GPU to an H100. Blocks must therefore be independent: no inter-block sync within a kernel (barring cooperative groups). This is a favorite "why is it designed this way" question.
CUDA ↔ SYCL Terminology (sidebar — interviews speak CUDA)
Interviews use CUDA vocabulary; SYCL/oneAPI is the portable equivalent (relevant mainly at Intel). The one mapping people get wrong: warp = sub-group, block = work-group — not "warp = work-group."
.x axis maps to SYCL's highest dim.
Pre-Volta, a warp shared one program counter — divergent threads couldn't progress independently. From CC 7.0 each thread has its own PC and can diverge/reconverge at sub-warp granularity. This breaks legacy warp-synchronous code; use __syncwarp() and the _sync primitive variants (Module 7).
Conceptual Round
1. Why does CUDA guarantee no execution order between thread blocks?
2. A block of 100 threads runs on an SM. How many warps?
• "warp = work-group" is wrong — warp = sub-group, block = work-group.
• A block runs entirely on one SM; it never spans two.
• Zero-overhead warp switching is the latency-hiding mechanism (Module 5).
• Post-Volta, don't rely on implicit warp-synchronous behavior.
The Memory Hierarchy
Order the levels by latency and bandwidth with rough numbers — that's a direct interview question. The non-obvious one: on modern GPUs, L1 and shared memory are the same physical SRAM.
The latency/bandwidth ordering, the scope and lifetime of each level, and that using more shared memory shrinks available L1 (they share SRAM). Rough HBM bandwidth of a current GPU is fair game.
Five Levels, Capacity vs Speed
The whole game is keeping data in fast on-chip storage and minimizing global traffic. A global load costs ~500 cycles — but with enough resident warps, the scheduler hides it entirely (Module 5).
Conceptual Round
1. Why is ~500-cycle global latency usually not a throughput problem?
2. You request 100 KB shared memory per block on an H100. Side effect?
• Forgetting shared mem and L1 are the same SRAM (a tunable carveout).
• Confusing scope: registers are per-thread, shared is per-block, L2/global are device-wide.
• Quoting a register/shared bandwidth as if it were HBM — they differ by ~10×.
• The point of the hierarchy: raise reuse so you touch HBM as little as possible (Modules 4, 6).
Memory Coalescing & Access Patterns
One of the most-asked CUDA questions. Define coalescing, know what breaks it, and explain why SoA beats AoS on a GPU. It's the single biggest "free" speedup in most kernels.
That coalescing is a warp property (which thread touches which address), not caching — and that you can reason about transaction count for a given pattern. AoS-vs-SoA is a classic follow-up.
128 Bytes at a Time
Global memory is served in fixed-size transactions. If a warp's 32 threads read 32 consecutive, aligned 4-byte words (128 bytes), the hardware coalesces them into the minimum number of transactions (100% efficiency). Any stride or misalignment wastes transactions:
i reads element base+i → four 32-byte sectors, ~100% efficiency.
Spot the Bug: Uncoalesced Access
This kernel updates .x for every particle and runs at ~3% of memory bandwidth. Click the design decision responsible.
struct Particle { float x, y, z; }; Particle p[N]; // AoSint i = blockIdx.x*blockDim.x + threadIdx.x;p[i].x += dt * v[i]; // warp strides by 12 bytesStore float x[N], y[N], z[N] separately. Now a warp's x[i] reads are contiguous → coalesced → ~100% efficiency. SoA is almost always the GPU-friendly layout.
Conceptual Round
1. Two kernels, same bytes and FLOPs, one 6× faster. Why?
2. Coalescing is fundamentally about...?
• Coalescing needs alignment, not just contiguity — a misaligned base still fragments.
• It's easy to break with a transpose or a bad / vs % in the index.
• AoS is natural in CPU code but poison on GPU — prefer SoA.
• Same guidance in SYCL: adjacent work-items → adjacent addresses.
Shared Memory & Bank Conflicts
A top separator question. Explain the 32-bank model, what causes a conflict, and the +1 padding fix. Bank conflicts happen within a warp — that phrasing matters.
The 32-bank model, why the +1 padding trick works, the broadcast exception, and that __syncthreads() is block-level (believing it syncs the whole grid is a red flag).
32 Banks, Serialized on Conflict
Shared memory is a programmer-managed scratchpad used to stage data (tiling) so global memory is read once and reused. It's split into 32 banks; successive 32-bit words map to successive banks. A bank conflict = N threads in a warp hit N different addresses in the same bank → the accesses serialize (N× slower).
If all threads read the same address, hardware broadcasts — no conflict. Conflicts only arise from different addresses in the same bank. A full-warp stride-32 access lands all 32 threads in one bank → a 32-way conflict.
The +1 Padding Trick
// conflict: column access, all in bank 0
__shared__ float tile[32][32];
float v = tile[threadIdx.x][0]; // 32-way!
// fix: pad the leading dimension
__shared__ float tile[32][33]; // +1
float v = tile[threadIdx.x][0]; // conflict-freei col 0 = element i*32 = bank (i*32)%32 = 0 → every row's col 0 is bank 0 → 32-way conflict.[33], each row shifts by one bank, so consecutive rows' col 0 land in different banks. Costs one word per row.Conceptual Round
1. __shared__ float s[32][32], threads read s[threadIdx.x][0]. Conflict degree + fix?
2. When is same-bank access NOT a conflict?
• __syncthreads() is block-level, not device-level — it does not sync the whole grid.
• Bank conflicts are a within-a-warp phenomenon tied to 32 banks / 4-byte words.
• You need __syncthreads between writing shared memory and reading what other threads wrote.
• Diagnose with Nsight's shared_ld_bank_conflict metric (Module 8).
Occupancy & Latency Hiding
The single best separator in GPU-performance interviews: higher occupancy is not always faster. Know the limiters, the Little's-Law intuition, and why forcing occupancy can backfire via register spilling.
Define occupancy, name the three limiters, and — the key signal — explain that once you have enough warps to hide latency, more occupancy gives nothing, and cutting registers to chase it can cause spilling.
Latency Hiding = Enough Warps in Flight
Occupancy = active warps per SM ÷ max possible. GPUs hide latency by keeping many warps resident — when one stalls, another issues. This is Little's Law: to keep the pipeline full you need concurrency = latency × throughput in-flight. At ~500-cycle latency and ~1 instr/cycle, you need hundreds of independent warp-instructions in flight.
Higher Occupancy Is NOT Always Faster
Once you have enough warps to hide latency, more occupancy buys nothing. Worse: forcing high occupancy by cutting registers-per-thread can cause register spilling to local memory (GMEM-backed, slow). Many high-ILP kernels run best at ~50% occupancy. The rule: low occupancy hurts latency hiding, but max occupancy ≠ max speed.
# SM: 65536 registers, max 2048 threads
# 32 regs/thread:
65536 / 32 = 2048 threads = 100%
# bump to 40 regs/thread:
65536 / 40 = 1638 -> ~1536 = 75%__launch_bounds__ / -maxrregcount to trade registers for occupancy deliberately — and measure, because more isn't always better.The Separator Round
1. 96 regs/thread, 256-thread blocks, 65536 regs/SM. Max resident blocks?
2. You raise occupancy 25% → 75% but the kernel gets slower. Cause?
• "Maximize occupancy always" is the wrong answer — it's a means to hide latency.
• Register spilling is silent; check -Xptxas -v for spill loads/stores.
• High-ILP kernels can beat high-occupancy ones (fewer warps, more work each).
• Use the Occupancy Calculator / Nsight to find which resource is the limiter.
The Roofline Model
The highest-frequency mental model across every GPU interview loop. Compute arithmetic intensity, find the ridge point, and classify a kernel memory- vs compute-bound — then you know what to optimize.
Whether you can compute the ridge point from hardware specs, classify a kernel, and say what to optimize. Bonus signal: knowing ridge points climb each GPU generation, so more kernels are memory-bound over time.
Arithmetic Intensity & the Ridge Point
Arithmetic intensity = FLOPs ÷ bytes moved. Attainable performance is:
The ridge point = peak compute ÷ peak bandwidth [FLOP/byte]. AI below it → memory-bound (optimize data movement). AI above it → compute-bound (optimize math / use tensor cores).
Ridge points are rising each generation — H100 BF16 tensor-core ridge ≈ 295 FLOP/byte, up from A100's ≈ 156 — so kernels are increasingly memory-bound on newer hardware.
Worked Example: SAXPY vs SGEMM
y = a*x + y: 2 FLOPs, 12 bytes (read x, read y, write y) → AI ≈ 0.17. On H100 (ridge 295) that's far left → max ~0.56 TFLOP/s of 989 peak. Only fewer bytes (fusion, lower precision) helps — never more compute.
2N³ FLOPs, ~16N² bytes → AI = N/8, grows with N. Large matmuls cross the ridge → compute-bound → map beautifully onto tensor cores. This is why matmul dominates ML.
The Estimation Round
1. Kernel AI = 10 FLOP/byte on H100 (ridge 295). Bound + action?
2. How do you increase a kernel's arithmetic intensity?
The roofline reappears in GEMM from Scratch (each optimization raises AI) and in Triton ("is my kernel memory- or compute-bound?"). Learn it once here; reuse it everywhere. It's arguably the highest-frequency thread across every GPU interview.
• "Peak" is precision-specific — FP32 vs FP16 tensor-core peak differ >10×. Always state precision.
• Small/skinny problems are latency- or launch-bound, not roofline-bound.
• AI is measured between specific memory levels — be clear which (HBM vs SMEM).
Warp-Level Ops, Divergence & Reductions
Warp divergence, shuffle reductions, and atomic contention all appear verbatim in question banks. The parallel reduction is GPU-interview canon — know the hierarchical warp→block→global pattern.
That divergence cost = serialization of branches within a warp (not "ifs are slow"), why _sync primitives exist post-Volta, and how to cut atomic contention with a hierarchical reduction.
Warp Divergence
A warp shares one instruction stream. If threads take different branches of a data-dependent conditional, the warp executes each path serially, masking off threads not on the current path. if (threadIdx.x % 2) splits every warp → both paths run → ~2× cost. Divergence is confined to a warp — different warps are free. Minimize by aligning branches to warp size (branch on threadIdx.x / 32, not per-thread).
// register-to-register, no shared mem, no sync
for (int off = 16; off > 0; off /= 2)
val += __shfl_down_sync(0xffffffff, val, off);
// lane 0 now holds the warp sum__syncthreads.0xffffffff) is required post-Volta (independent thread scheduling).The Hierarchical Reduction — Match Each Stage
Summing a big array is GPU-interview canon. The efficient pattern is a three-level hierarchy that minimizes atomic contention. Drag each technique to its stage.
__shfl_down_syncshared memory + __syncthreadsone atomicAdd per blockatomicAdd per thread to one addressConceptual Round
1. if (threadIdx.x % 2 == 0) A(); else B(); — cost?
2. Why is a __shfl_down_sync reduction faster than shared memory for the final 32 elements?
• Legacy __shfl/__ballot (no mask) are deprecated — use _sync variants (independent thread scheduling). Adding __syncwarp before a legacy call does not make it safe.
• A single-address atomicAdd from thousands of threads serializes — reduce hierarchically first.
• Divergence across warps is free; only within a warp costs.
Profiling, Concurrency & the Optimization Workflow
Tie it together: measure → find the dominant bottleneck → fix one thing → re-measure. Plus streams/overlap — a high-frequency topic. Framed as methodology you narrate, not tool-clicks.
Your measure-first workflow, reading a profiler to classify memory- vs compute- vs latency-bound, and using streams to overlap compute with transfers. "How would you find the bottleneck?" is the real question — never guess.
The Systematic Loop
Nsight Compute's SOL shows % of peak for Compute vs Memory. The higher one tells you which roof you're near; both low → latency-bound.
Check coalescing/sector efficiency, reduce bytes (fusion, precision), stage in shared memory, exploit L2.
Reduce instructions, use tensor cores, fix divergence, faster math intrinsics.
Low both-SOL + high long-scoreboard stalls → too few eligible warps. Then re-measure.
Streams: Overlap Compute & Transfers
A high-frequency topic: host↔device PCIe (~16–64 GB/s) is far slower than device bandwidth (~3 TB/s), so you overlap transfers with compute using streams and pinned memory.
// copy chunk i+1 while computing chunk i
cudaMemcpyAsync(d+off, h+off, sz,
cudaMemcpyHostToDevice, stream[i]);
kernel<<<g, b, 0, stream[i]>>>(d+off);
cudaMemcpyAsync(h+off, d+off, sz,
cudaMemcpyDeviceToHost, stream[i]);cudaMallocHost).The Diagnosis Round
1. Nsight: Memory SOL 85%, Compute SOL 20%, dominant stall "long scoreboard." Diagnosis?
2. Both Compute and Memory SOL ~25% yet the kernel is slow. What's wrong?
Execution model → memory hierarchy → coalescing → bank conflicts → occupancy → roofline → warps & reductions → profiling. The through-line every interviewer rewards: classify the bound first (memory / compute / latency), then optimize the thing that's actually limiting you — measured, not guessed.
• Uncoalesced/AoS access; unnecessary global round-trips; bank conflicts; single-address atomic contention.
• Chasing occupancy into register spills; optimizing FLOPs on a memory-bound kernel.
• Unhidden PCIe transfers — overlap them with streams.
• Apply this next in GEMM from Scratch, where each step is a measured march up the roofline.