CUDA, Triton & Kernel Engineering
30 questionsDONEUNLOCKEDLOCKED
CUDA & Triton Kernel Interview Questions
Coalescing, shared memory and bank conflicts, occupancy, fusion, tiled GEMM, FlashAttention internals, Triton, CUTLASS, Nsight profiling and torch.compile: the live-coding and take-home round at NVIDIA, Fireworks, Together and the labs' performance teams.
Grounded in real AI infrastructure interview loops and written to a senior-engineer editorial bar, with every number worked and every diagram hand-built.
You have 10 free answers unlocked here.Sign in free for 10 more · 10 are premium.
01–12Foundationsthe vocabulary every loop assumes you already have0/12 done
13–23Core loopsthe questions every loop actually asks0/11 done
24–30Field scenariosthe messy, half-specified problems from real deployments0/7 done
The concepts behind CUDA, Triton & Kernel Engineering
The vocabulary and mental models these questions assume, from our curriculum. Start with the foundations free; the deeper, interview-defining ideas are part of premium.
Foundational
CUDA Programming ModelCUDA splits a program into a host that allocates, copies and enqueues work, and a device that runs thousands of identical threads organized as a grid of blocks. Getting the split right, and knowing that a launch returns before the kernel runs, decides whether your first live-coding kernel produces a correct number or a silent zero.Core
Memory CoalescingA warp's 32 threads issue one memory request together, and the hardware serves it in 32-byte sectors. Coalescing is arranging addresses so those sectors are full of bytes the warp will use. It decides whether a bandwidth-bound kernel moves at the HBM rate or at an eighth of it, and it is the pattern NVIDIA's trace-classification interview question tests.Sign in
Advanced
Shared Memory and Bank ConflictsShared memory is the programmer-managed SRAM inside each SM, split into 32 four-byte banks that serve one word each per cycle. When several lanes of a warp hit the same bank at different addresses the access serializes, and a 32-way conflict makes a shared-memory-bound loop run over ten times slower. Padding, XOR swizzles, cp.async and TMA are the tools that decide whether a tiled kernel gets the bandwidth it staged data for.🔒 Premium
Advanced
Occupancy and Register PressureOccupancy is the fraction of an SM's 64 warp slots that are resident, and it is capped by the 65,536 registers and 228 KB of shared memory each block consumes. It decides how much memory latency the hardware can hide for free, but the fastest kernels on a GPU routinely run at 25 percent, so the interview skill is knowing when to raise it and when to stop.🔒 Premium
Foundational
Kernel FusionAn elementwise or reduction kernel does a few FLOPs per byte and runs at HBM speed, so a chain of five of them costs five trips through HBM for work that needs one. Fusion collapses the chain into a single kernel that keeps intermediates in registers. It is the first lever for anything memory-bound, and knowing what it cannot fix (weight reads in decode, the GEMMs themselves) is what the interview is really testing.Advanced
Tiled Matrix MultiplicationA matrix multiply has enough reuse to be compute-bound, but only if the kernel captures that reuse in shared memory and registers instead of re-reading HBM. Tiling is how: a block owns an output tile, streams K-slices of A and B through shared memory, and each thread accumulates a small register tile. It is the live-coding exercise that separates people who know the roofline from people who have climbed it.🔒 Premium
Advanced
FlashAttention InternalsStandard attention writes the N x N score matrix to HBM and reads it back, which makes it memory-bound and quadratic in memory. FlashAttention tiles Q, K and V through shared memory, keeps a running max and sum so the softmax never needs the full row, and recomputes scores in the backward pass. Knowing the online-softmax rescale, why FlashAttention-2 flipped the loop order, and what FlashAttention-3 overlaps on Hopper is the difference between naming the paper and being able to write the kernel.🔒 Premium
Advanced
CUTLASS and Tensor Core KernelsCUTLASS is NVIDIA's template library for building GEMM-shaped kernels that run tensor cores at near cuBLAS speed while letting you change the data types, the tile shapes and the epilogue. Its hierarchy (device, kernel, collective mainloop, tile, instruction) is the vocabulary of every tensor-core discussion, and knowing when it beats calling cuBLAS or writing Triton is the judgment question kernel interviews end on.🔒 Premium
