The question that sorts candidates on an AI-infra loop is rarely write a matrix multiply. It’s some version of here’s a kernel running at 15% of peak memory bandwidth, tell me why and make it faster. Nvidia asks it. So do Modal, Together, and Fireworks, and the chip startups like Groq and Cerebras, each with their own accent. What they’re checking is whether you think about the GPU as a memory system with thousands of threads attached, or whether you memorized the runtime API and hoped for the best.
If you can talk fluently about warps, coalescing, and occupancy, and you profile before you guess, you’re already past most of the field. Here’s what those rounds actually probe and how a strong answer sounds out loud.
Threads run in warps, and the warp is where the money is
A warp is 32 threads that execute the same instruction in lockstep. That single fact explains most GPU performance questions. A common opener: what happens when half the threads in a warp take the if branch and half take the else? The answer interviewers want is that both paths run serially with the inactive threads masked off, so a branchy kernel can quietly cost you 2x or worse. This is warp divergence, and the good follow-up from you is that it only matters when threads within the same warp diverge. If the branch splits cleanly on warpId, there’s no penalty, because every thread in a given warp agrees.
The reason this comes first is that it sets up everything else. Threads aren’t independent little CPUs. They move in packs of 32, and the hardware wants those 32 to do the same thing and touch nearby memory at the same time.
Coalescing is the answer to most “why is this slow” questions
When the 32 threads in a warp read 32 consecutive addresses, the memory controller serves them in one or a few wide transactions. When they read scattered or strided addresses, that single logical load fans out into many transactions and you throw away most of your bandwidth. Coalescing is the name for the good case, and a huge share of “make this kernel faster” answers reduce to “make the global accesses coalesced.”
The canonical question is the matrix transpose. Read row by row and the loads are coalesced. Write column by column and the stores are strided by the matrix width, so the writes are uncoalesced and the kernel crawls. Interviewers want you to reach for shared memory: load a tile of the input with coalesced reads, transpose inside shared memory, then write the tile back out with coalesced writes. The extra detail that earns points is padding the shared array by one column (tile[32][33]) so successive rows land in different banks and you avoid shared-memory bank conflicts. Bringing that up unprompted is a strong signal, because it means you’ve actually profiled a transpose and seen the bank-conflict stall.
Knowing where data lives is half of this. The memory hierarchy is worth having in your head cold, because the whole optimization game is moving the hot data up the table and keeping it there.
| GPU memory space | Visible to | Managed by | Approx. access latency | What you use it for |
|---|---|---|---|---|
| Registers | A single thread | Compiler | ~1 cycle | Per-thread scalars; fastest storage, but the register budget caps how many warps fit |
| Shared memory | All threads in a block | You, explicitly | ~20-30 cycles | Tiling, reductions, and staging data to turn uncoalesced global access into coalesced access |
| L2 cache | The whole device | Hardware | ~200 cycles | Automatic reuse across blocks; you influence it only through access patterns |
| Global memory (HBM) | Whole device and host | You | ~400-600 cycles | Main inputs and outputs, model weights, the KV cache |
| Constant memory | Whole device, read-only | You | ~1 cycle on a cache hit | Small parameters broadcast identically to every thread |
| Local memory | A single thread | Compiler | Global-memory latency | Register spills; seeing it in the profiler is a warning sign, not a plan |
Occupancy is a lever, not a goal
Somewhere in the loop you get: this kernel runs at 50% occupancy, how do you get it to 100%? It’s a trap, and chasing 100% is the wrong instinct. Occupancy is the ratio of active warps to the maximum the hardware can hold, and its only job is to hide latency: while one warp waits on a 500-cycle global load, the scheduler runs another. Past the point where you have enough warps to cover that latency, more occupancy buys nothing, and pushing for it usually costs registers or shared memory per thread, which can make each thread slower.
The answer that lands is that you’d raise occupancy only if the profiler shows the kernel is latency-bound with too few warps in flight, and that a kernel can hit peak throughput at 50% occupancy if each thread does enough independent work to keep the pipeline full. If you know the reference, name Vasily Volkov’s “Better Performance at Lower Occupancy,” which made exactly this point using instruction-level parallelism. Interviewers at Nvidia in particular light up when you treat occupancy as one input to latency hiding rather than a score to maximize.
Sum a billion floats without wasting the machine
Parallel reduction is the most-asked coding question in this space because it forces you to think about all of it at once: coalesced loads, avoiding divergence, using shared memory or warp primitives, and not letting a serial tail dominate. The naive version has every other thread idle after the first step and hammers shared memory. The version they want uses a grid-stride loop for the loads and then reduces inside each warp with __shfl_down_sync, which trades shared memory for register-to-register shuffles across a warp.
__global__ void reduce(const float* in, float* out, int n) {
float sum = 0.0f;
for (int i = blockIdx.x * blockDim.x + threadIdx.x;
i < n; i += gridDim.x * blockDim.x)
sum += in[i]; // grid-stride, fully coalesced loads
// reduce the 32 lanes of this warp with no shared memory
for (int offset = 16; offset > 0; offset >>= 1)
sum += __shfl_down_sync(0xffffffff, sum, offset);
if ((threadIdx.x & 31) == 0) // lane 0 of each warp
atomicAdd(out, sum); // one atomic per warp, not per thread
}The point isn’t to reproduce this from memory. It’s to explain why each choice exists: the grid-stride loop keeps loads coalesced regardless of grid size, the shuffle reduction skips shared memory and its synchronization, and doing one atomicAdd per warp instead of per thread cuts the atomic traffic by 32x. If you can also say when you’d switch to a two-pass reduction or CUB’s DeviceReduce instead of hand-rolling, you sound like someone who ships kernels rather than someone who read a blog post.
Where kernels meet serving
AI-infra roles push past raw CUDA into why kernels get written the way they do for models. Kernel fusion is the recurring theme: every separate kernel means another round trip to HBM, so fusing an elementwise op into the matmul that feeds it saves a full read and write of a large tensor. FlashAttention is the interview-favorite example, and a good candidate can explain that it never materializes the full N-by-N attention matrix in global memory. It tiles the computation to keep the working set in shared memory and recomputes what it needs, trading a bit of extra math for a large cut in memory traffic.
That trade only makes sense if you can tell whether a kernel is memory-bound or compute-bound, which is the roofline question. A sharp version asked at inference shops: during single-request decoding, is attention memory-bound or compute-bound? The answer is memory-bound. At batch size one you re-read the weights and the growing KV cache for every token while doing very little arithmetic per byte, so arithmetic intensity is low and you sit under the memory roof. That single insight is why batching, paged KV cache, and quantization exist, and why FP8 or INT8 on tensor cores helps: smaller data types move fewer bytes and feed the tensor cores faster once you’re compute-bound.
The same questions, different accents
Nvidia goes deepest on architecture. Expect Nsight Compute profiler screenshots to reason about, questions that reach down to PTX or SASS, and a real interest in whether you understand tensor cores and the memory subsystem rather than just calling cuBLAS. Groq and Cerebras test a different reflex, because their hardware doesn’t look like a GPU. Groq’s LPU is a deterministic, compiler-scheduled dataflow machine with no caches to hide behind, so their questions lean on static scheduling and knowing exactly where every byte is at every cycle. Cerebras runs wafer-scale with weight streaming, and they want people comfortable with model and pipeline parallelism at a scale where the interconnect is the bottleneck.
The inference startups, Modal, Together, Fireworks, and Baseten among them, blend CUDA with systems. You’ll still get a coalescing or occupancy question, but it sits next to your inference pod is getting OOMKilled on an 80GB GPU, walk me through it. The expected answer connects the KV cache growing with sequence length and batch size to the fix: paged attention, a smaller max batch, quantizing the cache, or spilling to host. That crossover from kernel to cluster is the whole job at those companies, and the interview mirrors it.
The strongest signal across all of them is small and easy to show: reach for the profiler before you theorize. A candidate who says “I’d look at the achieved memory throughput and the warp stall reasons in Nsight before changing a line” reads as someone who has debugged a slow kernel for real, which is worth more than any single trick you can recite.
