CUDA Interview Questions for ML Engineers

What CUDA interview questions do ML engineers get asked?

CUDA interview questions for ML engineers cover the GPU execution model, the memory hierarchy, kernel performance and classic kernels such as reduction, transpose, tiled matrix multiply and softmax. ML infrastructure roles add questions on kernel fusion, FlashAttention, LLM decode and KV cache sizing. Interviewers mostly want to hear you reason about how much memory each kernel moves.

The questions below come from what candidates commonly report. Each answer lists what a strong reply covers and the follow-up that tends to come next, with a TensorTonic problem where one matches.

Which GPU execution model questions come up first?

Interviews usually open with how threads are organised and scheduled, because every later answer depends on it.

How do threads, blocks and grids map onto the data?

A launch creates a grid of blocks, and each block of up to 1,024 threads runs on one streaming multiprocessor (SM) and shares its shared memory. Each thread computes blockIdx.x * blockDim.x + threadIdx.x and checks it against the array length, because the grid is rounded up to whole blocks. The follow-up asks how to cover an array larger than the grid, which is what a grid-stride loop does.

Try Vector Addition as a CUDA Kernel.

What is a warp, and what is warp divergence?

A warp is 32 threads that the SM issues one instruction to at a time. When lanes of one warp take different sides of a branch, each side runs in turn with the other lanes masked off, while a branch that splits on warp boundaries costs nothing extra. Strong candidates add that since the Volta generation, threads have independent program counters, so warp-level code should use __syncwarp and the _sync shuffle intrinsics.

Try Implement Leaky ReLU as a CUDA Kernel.

What is the difference between __syncthreads and cudaDeviceSynchronize?

__syncthreads is a barrier for the threads of one block inside a kernel, and cudaDeviceSynchronize makes the host wait for all launched GPU work. Blocks have no barrier between them inside an ordinary kernel, so they combine results through global memory with atomics or a second launch. The trap is a __syncthreads inside a branch some threads skip, which is undefined behaviour and often hangs.

What memory questions do CUDA interviews ask?

Memory questions take up the largest share of most CUDA interviews, because most ML kernels spend their time waiting on memory.

What is the GPU memory hierarchy?

From fastest to slowest, a GPU has per-thread registers, shared memory and L1 inside each SM, an L2 cache shared by the chip, and global memory in HBM or GDDR. Latency grows from about a cycle to several hundred, so quote orders of magnitude rather than exact figures. The follow-up is usually local memory, which holds register spills and lives in global memory despite its name.

What is memory coalescing?

Coalescing is the hardware merging one warp's global loads into as few transactions as possible. When 32 threads read 32 consecutive floats, the warp touches one 128-byte span, while reading down a column of a row-major matrix can cost a separate 32-byte sector per thread. The follow-up is a transpose with both reads and writes coalesced, which needs a shared memory tile.

Try Global-Memory Coalescing Counter.

What is a shared memory bank conflict, and how do you fix it?

Shared memory has 32 banks of 4-byte words, and when two lanes want different words from the same bank, the bank serves them one after the other. Reading a column of float tile[32][32] puts every element in one bank and takes 32 passes. Padding to [32][33] shifts each row by one bank and removes the conflict. Switch patterns below to see each case.

__shared__ float tile[32][32];
float v = tile[lane][3];

Try Shared-Memory Bank Conflict Analyzer.

How do interviewers test kernel performance?

Interviewers usually show you a slow kernel and ask what limits it, to see whether you can name the bottleneck before changing code.

What is occupancy, and what limits it?

Occupancy is resident warps on an SM divided by the most it can hold, and registers per thread, shared memory per block and block size all cap it. It exists to hide memory latency, so a kernel that keeps more data in registers can win at lower occupancy.

Try GPU Occupancy Calculator.

A kernel reaches a fraction of peak bandwidth. How do you find out why?

You profile it first. A timeline profiler such as Nsight Systems shows whether the GPU sits idle between launches and copies, and Nsight Compute reports one kernel's achieved bandwidth and why its warps stall. The usual suspects are uncoalesced access, bank conflicts, register spills and too little work per thread.

Try Measure Kernel Timeline Utilization.

Is this kernel memory-bound or compute-bound?

The roofline model answers it with arithmetic intensity, the FLOPs a kernel does per byte it moves. The ridge point is peak FLOP/s divided by peak bandwidth, in the low hundreds of FLOPs per byte for half-precision math on recent data-centre GPUs. FP32 vector addition does one add per 12 bytes, about 0.08 FLOPs per byte, so it is memory-bound everywhere. A 4,096 by 4,096 half-precision matrix multiply does about 137 billion FLOPs over at least 101 MB, roughly 1,365 FLOPs per byte, which puts it well past the ridge.

How do streams and pinned memory overlap copies with compute?

Work issued to different non-default streams can run at the same time, but cudaMemcpyAsync only overlaps with compute when the host buffer is pinned, for example with cudaMallocHost. The follow-up asks how to hide transfers in a pipeline, and the answer is to split the batch across a few streams so one chunk copies while another computes.

Which CUDA kernels are you asked to write in an interview?

Candidates commonly report the same short list of kernels, and each one tests a specific idea.

A strong reduction answer looks roughly like this, with a grid-stride loop, a warp shuffle sum, one shared memory slot per warp and one atomic per block. It assumes the block size is a multiple of 32.

__device__ float warp_sum(float v) {
    for (int offset = 16; offset > 0; offset >>= 1)
        v += __shfl_down_sync(0xffffffff, v, offset);
    return v;
}

__global__ void sum_kernel(const float* x, float* out, int n) {
    __shared__ float partial[32];
    float v = 0.0f;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += gridDim.x * blockDim.x)
        v += x[i];
    v = warp_sum(v);
    int lane = threadIdx.x % 32;
    int warp = threadIdx.x / 32;
    if (lane == 0) partial[warp] = v;
    __syncthreads();
    if (warp == 0) {
        v = lane < blockDim.x / 32 ? partial[lane] : 0.0f;
        v = warp_sum(v);
        if (lane == 0) atomicAdd(out, v);
    }
}

Expect to be asked who zeroes the output before launch, and why the atomics make the last bits of a float sum depend on the order in which blocks finish.

What GPU questions are specific to ML infrastructure roles?

ML infrastructure interviews ask how these ideas show up in attention, LLM serving and multi-GPU training.

Why does kernel fusion speed things up?

Each separate elementwise kernel reads its input from global memory and writes its output back, so a chain of bias add, activation and dropout moves one tensor several times. Fusing them reads and writes it once, which raises arithmetic intensity for memory-bound operations and saves launch overhead.

Try Fused Matmul + Bias + ReLU in Triton.

Why is FlashAttention faster than standard attention?

Standard attention writes the N by N score matrix to HBM and reads it back. FlashAttention loads tiles of Q, K and V into on-chip SRAM, keeps a running max and sum per row with an online softmax, and never stores the full matrix. It does about the same FLOPs and moves far fewer bytes. The FlashAttention lesson shows the tiling step by step.

Try FlashAttention Forward Kernel in Triton.

Why is LLM decode memory-bound when prefill is compute-bound?

Prefill reuses each weight across every prompt token, so its matrix multiplies sit right of the ridge. Decode makes one token per sequence per step, so each step reads every weight and the KV cache for about two FLOPs per weight, around one FLOP per byte in half precision at batch size one. Batching more sequences raises that intensity.

How much memory does the KV cache take?

The KV cache takes 2 × layers × KV heads × head dimension × sequence length × batch size × bytes per element. With 32 layers, 8 KV heads, a head dimension of 128 and 8,192 tokens in FP16, that is exactly 1 GiB per sequence. To fit more, the usual levers are grouped-query attention, a quantized cache and paged allocation.

Try Plan KV-Cache Capacity.

What is all-reduce?

All-reduce sums a buffer across GPUs and leaves the result on each of them, which is how data-parallel training averages gradients. In a ring across N GPUs each one sends about 2(N - 1)/N times the buffer, so the slowest link sets the pace.

Try Calculate Collective Effective Bandwidth.

How should you answer a CUDA question you are not sure about?

Reason from bytes and FLOPs, which gets you to a sensible answer without the exact fact. If you are asked why a softmax kernel is slow, you can say it does a few FLOPs per element and should be memory-bound, that you would compare its achieved bandwidth with peak in a profiler, and that you would fuse its three passes over each row into one. When the answer depends on the GPU generation, say so instead of guessing a number.

How do you practise for a CUDA interview?

You practise by writing kernels that run, pass tests and get timed. The CUDA practice problems guide orders the 35 problems in the CUDA Kernels plan, which run on real cloud GPUs. The Triton plan, the LLM Internals lessons and the FlashAttention project cover the ML side. For timed practice on the ML side, the LLM Internals mock on assessments is a 45-minute, three-problem Pro round on attention and KV-cache code, and the ML engineer interview guide covers the rest of the loop.

Related guides

• CUDA practice problems: from vector add to attention
• TensorTonic or LeetGPU: where should you learn CUDA?
• How to prepare for an ML engineer interview
• ML interview mistakes: the five that cost offers
• ML coding interview questions: 20 to implement
• ML mock interview: how to check you are ready