CUDA Interview Questions and Answers (NVIDIA and GPU Roles, 2026)

The short answer: interviews for CUDA and GPU roles, at NVIDIA and at the many companies that build on its GPUs, test four things: the execution model (threads, blocks, grids, warps and SMs), the memory hierarchy and how to use it well (coalescing, shared memory, bank conflicts), concurrency (streams and overlapping transfers with compute), and whether you can write and optimize a kernel, usually a reduction, a matrix multiply or a transpose. Below are the questions with short answers checked against NVIDIA’s CUDA Programming Guide, then the coding exercises and the inference questions that ML roles add.

Execution model

What is a kernel, and how is it launched?

A kernel is a function that runs on the GPU (the device), launched from the CPU (the host). The launch specifies an execution configuration: the grid and thread block dimensions, plus optional settings such as the stream. Every thread runs the same kernel code on different data.

How are threads organized?

Threads are grouped into thread blocks, and blocks into a grid. Blocks and grids can have one, two or three dimensions. Each thread finds its position with built-in variables (threadIdx, blockIdx, blockDim), for example a global index blockIdx.x * blockDim.x + threadIdx.x. GPUs with compute capability 9.0 and higher add an optional level between blocks and the grid: thread block clusters.

What is a warp?

Within a thread block, threads are organized into groups of 32 called warps. The warp is the unit the hardware schedules, which is why warp size shows up in coalescing, bank conflicts and divergence.

What is warp divergence, and why does it matter?

When threads in the same warp take different branches, the paths are executed with the inactive threads masked off. Utilization is highest when all threads in a warp follow the same control flow, so avoid data-dependent branches that split warps, or arrange data so that whole warps take the same path.

What is a streaming multiprocessor (SM)?

The GPU is a collection of SMs, grouped into Graphics Processing Clusters. Each SM has its own register file and shared memory. When a kernel launches, the scheduler assigns thread blocks to SMs as resources allow.

Can you rely on the order in which blocks run?

No. The scheduler makes no ordering guarantees, and the application can’t control which SM runs which block. Blocks must be independent; correctness can’t depend on a scheduling order.

What is occupancy?

How many warps are active on an SM compared with the maximum it supports. It’s limited by each block’s use of registers, shared memory and threads. Higher occupancy helps hide memory latency, but it’s a means, not the goal: a kernel can be fast at modest occupancy if it has enough independent work per thread.

Memory

Describe the CUDA memory hierarchy.

• Registers: private to each thread, fastest.
• Shared memory: on-chip, shared by all threads in a block; used to reuse data and cooperate.
• Local memory: private to a thread but stored in device memory, so it’s as slow as global memory; used for register spills and some arrays.
• Global memory: the GPU’s DRAM, visible to all SMs; large and slow compared with on-chip memory.
• Constant memory: read-only for kernels, declared with __constant__, cached, and best when all threads read the same value.

What is memory coalescing?

When the threads of a warp access consecutive addresses, their global memory accesses are combined into as few transactions as possible. Strided or scattered access splits them into many transactions and wastes bandwidth. Layouts that let consecutive threads read consecutive elements (structure of arrays rather than array of structures, for example) help.

What are shared memory bank conflicts?

Shared memory has 32 banks, and successive 32-bit words map to successive banks. When threads in the same warp access different addresses in the same bank, the accesses are serialized. The exception is several threads reading the same word, which is broadcast. A common fix is padding a 2D shared array by one column, as in a matrix transpose.

When do you need __syncthreads()?

It’s a barrier for all threads in a block. Use it between writing shared memory and reading what other threads wrote, for example after loading a tile. It must be reached by all threads in the block, so don’t put it inside a branch that only some threads take.

How do you handle several threads updating the same value?

With atomic operations (atomicAdd and others) or, better, by reducing contention first: each block accumulates in shared memory, then one atomic per block updates global memory.

Concurrency and transfers

What is a CUDA stream?

A sequence of operations, such as copies and kernel launches, that execute in order. Work launched without a stream goes to the default stream. Operations in different streams can run concurrently.

How do you overlap data transfers with computation?

Split the work into chunks and use several streams, so one chunk is copied while another is computed. For copies involving CPU memory to be asynchronous, the host buffers must be pinned (page-locked).

How do you find where a kernel is slow?

Measure before optimizing: Nsight Systems for the timeline (transfers, launches, gaps), Nsight Compute for a single kernel (memory throughput, occupancy, stalls). Then decide whether the kernel is memory-bound or compute-bound, because the fixes differ.

Coding exercises

Expect to write at least one kernel. The classics:

• Vector addition: the warm-up. Compute the global index and check it against the array length.
• Reduction (sum an array): load into shared memory, reduce in a loop with __syncthreads() between steps, and discuss improvements: sequential addressing to avoid divergence and bank conflicts, and warp-level shuffles for the last steps.
• Matrix multiplication: a naive version, then a tiled version that loads tiles into shared memory to reuse data.
• Matrix transpose: coalesced reads and writes through a shared memory tile, padded to avoid bank conflicts.
• Histogram: atomics, then per-block histograms in shared memory to reduce contention.
• Prefix sum (scan): harder; know the idea of a work-efficient scan even if you don’t write it in full.

A minimal kernel to be able to write without thinking:

__global__ void add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

// Launch with enough blocks to cover n elements
int threads = 256;
int blocks = (n + threads - 1) / threads;
add<<<blocks, threads>>>(d_a, d_b, d_c, n);

Talk through each choice: why the bounds check, why 256 threads per block (a multiple of the 32-thread warp), and what you’d measure next.

Inference and ML systems questions

Roles on inference and deep learning frameworks add questions like these:

• Why is inference often memory-bound? For small batches, especially in LLM decoding, each weight is read from memory for relatively little computation, so memory bandwidth, not arithmetic, sets the speed. Batching and quantization help.
• What does quantization trade off? Lower precision (such as INT8 or FP8) cuts memory traffic and raises throughput, at some risk to accuracy that has to be measured.
• What is a KV cache? In autoregressive transformers, the keys and values of past tokens are stored so each new token doesn’t recompute them. It saves compute but uses a lot of memory as sequences grow.
• What is kernel fusion? Combining several operations into one kernel so intermediate results stay on-chip instead of going through global memory.
• What are Tensor Cores for? Specialized units for matrix multiply-accumulate at reduced precision, used heavily in training and inference.

Behavioral questions

GPU roles still have behavioral rounds: a performance problem you tracked down, a design trade-off you made, a disagreement about an approach. Prepare two or three stories with numbers (“the kernel went from 12 ms to 3 ms because…”) and tell them with the STAR method. For the coding-round format, see our CoderPad guide and AI coding interviews.

Frequently asked questions

Sources

• NVIDIA, CUDA Programming Guide: Programming Model, Writing SIMT Kernels and Asynchronous Execution. Read on 29 September 2026.

NVIDIA and CUDA are trademarks of NVIDIA Corporation. Huru is not affiliated with or endorsed by NVIDIA.