Interview Bank 01 — GPU & Systems Deep-Dive (Phases 01–05)

Derive, don't recite. Each answer reasons from a principle. Full answers in the phases' WARMUP Interview Q&A sections; this bank is the index + extra drills.

Architecture (Phase 01)

Q: A100 has 1,555 GB/s HBM and 312 TFLOPS BF16. Ridge point? Implication for decode? Ridge = 312e12 / 2.0e12 ≈ 156 FLOP/byte. Batch-1 decode is AI ≈ 2 → ~80× bandwidth-bound; ceiling for a 14 GB FP16 model ≈ 2.0 TB/s ÷ 14 GB ≈ 140 tok/s. Serving economics = raise effective AI (batching, quantization). [P01 WARMUP Q3]

Q: Why no branch predictors / OoO on a GPU? Latency-hiding by oversubscription replaces latency-removal by prediction; those transistors become ALUs. [P01 WARMUP Q1]

Drill: warp executes if (threadIdx.x % 4 < 2) — cost? (~2×, two masked serialized paths; the % 4 < 2 splits each warp in halves.)

Drill: when is 100% occupancy the wrong goal? (Latency already hidden; cutting registers to raise occupancy spills — Volkov. [P01 WARMUP Q5])

CUDA (Phase 02)

Q: Kernel gets 8% of peak bandwidth — first three hypotheses? Uncoalesced access (sectors/request ≈ 32), latency-bound (too few warps), or the kernel isn't the bottleneck (nsys timeline). [P02 WARMUP Q2]

Q: Launch config for a 1920×1080 filter. 16×16 blocks, 120×68 grid, guard x<1920 && y<1080, shared-mem halo tile. [P02 WARMUP Q1]

Drill: what breaks with __syncthreads() in divergent code? (Deadlock — all block threads must hit the same barrier. [P02 WARMUP Q4])

Drill: a PR claims 3× speedup — what do you require before merge? (Reference test + sanitizer, honest timing, ncu limiter evidence, arch notes. [P02 WARMUP Q6])

Compilers (Phase 03)

Q: What happens compiling+running CUDA on a GPU that didn't exist at compile time? No matching cubin → driver JITs the embedded PTX → SASS at load, caches it; slow first load (the cold-start incident); fails if no PTX embedded. [P03 WARMUP Q1]

Q: Why does PTX exist? Virtual ISA for stability across GPU generations whose hardware ISAs change; splits slow AOT optimization from fast load-time lowering; CUDA's compatibility moat. [P03 WARMUP Q2]

Drill: how do you spot register pressure in SASS? (LDL/STL spills. [P03 WARMUP Q3])

Drill: why did Triton succeed? (Block-level abstraction — automates tiling/ staging, ~90% of hand-tuned, Python+MLIR. [P03 WARMUP Q4])

Drivers & containers (Phase 04)

Q: What happens at the OS level on cudaLaunchKernel? Init: open /dev/nvidia*, ioctls create context, mmap ring+doorbell, JIT PTX. Launch: write ring + tap doorbell — zero syscalls; GPU writes a fence; sync waits on it. [P04 WARMUP Q1]

Q: Why doesn't CUDA_VISIBLE_DEVICES provide isolation? Env var read by the untrusted process; enforcement is the devices cgroup, and even that gates whole devices — real isolation is MIG/dedicated. [P04 WARMUP Q2]

Drill: container sees the GPU but cudaInit fails — debug. (nvidia-smi on host; ldconfig for libcuda version skew; cgroup perms; strace the failing ioctl. [P04 WARMUP Q3])

Drill: why is there no GPU cgroup controller, and what fills the gap? (The KMD/ GPU owns GPU time, opaque to the kernel scheduler; MPS/MIG/time-slicing fill it. [P04 WARMUP Q2 / P06])

Runtime (Phase 05)

Q: nvidia-smi shows 78 GB used, framework says 51 GB — who's lying? Neither: reserved (driver's view, includes cache) vs allocated (live tensors). Gap is cache, not a leak. Watch largest-free-block for fragmentation. [P05 WARMUP Q2]

Q: Serving process OOMs after 3 days at steady load — hypotheses? External fragmentation (largest-free-block shrinking at <70% allocated), genuine leak (allocated trending up), pinned exhaustion, KV admission. [P05 WARMUP Q3]

Drill: design stream-ordered deallocation — what bug does it prevent? (Cross-stream use-after-free; free becomes an event in stream order. [P05 WARMUP Q4])

Drill: when do CUDA Graphs help and what can't they capture? (Many short kernels per iteration; not dynamic shapes/control flow/syncs. [P05 WARMUP Q5])

The cross-cutting synthesis question

Q: A 7B model serves at 30 tok/s on an H100 at batch 1. Walk me from silicon to the fix. Decode is bandwidth-bound (P01 roofline); 14 GB weights ÷ 3.35 TB/s ≈ 240 tok/s ceiling, so 30 means we're far below even that — check: is it actually compute or a pipeline stall (nsys, P02/P10), is the allocator thrashing (P05), is it launch- overhead-bound (CUDA Graphs, P05)? Then the economic fix is raising AI: batch (P07), quantize (P07). The answer threads P01→P05→P07 — exactly the role's breadth.