GPU Kernel

A GPU kernel is a function written in a serial style that the GPU executes simultaneously across many threads. In CUDA, a kernel is marked with the __global__ specifier and launched with a grid configuration; each of the thousands of threads runs the same code on a different slice of data. The kernel is the fundamental unit of work a GPU performs, and the efficiency of running AI models locally hinges on how well these kernels are written and scheduled.

The execution model: grids, blocks, and warps

When a kernel launches, its threads are organized into a grid of thread blocks, and each block holds a number of threads that can cooperate through fast shared memory on the same streaming multiprocessor. On NVIDIA hardware, threads are scheduled in groups of 32 called warps. All threads in a warp execute the same instruction at once on different data, a design called SIMT (Single Instruction, Multiple Thread). This is what lets a single matrix multiply spread across an entire GPU in one launch. When threads in a warp take different branches, the hardware must run each path in turn, so divergent code quietly wastes lanes — one reason well-written kernels avoid data-dependent branching in their innermost loops.

Why kernels matter for local inference

The speed of running an open-weight model at home is largely a story about kernels. Projects hand-tune kernels for attention, quantized matrix multiplication, and memory layout because a poorly written kernel can leave most of the GPU idle waiting on memory. Tools such as FlashAttention exist precisely because a smarter kernel can cut memory traffic dramatically without changing the underlying math — it fuses several operations into one pass so intermediate results never leave fast memory. When a new quantization format ships a fast kernel, that format becomes practical for home hardware; without one, the same format is a curiosity that looks good on paper and runs slowly in reality.

Occupancy and memory bound versus compute bound

Two ideas explain most kernel performance puzzles. Occupancy is how many warps a multiprocessor keeps in flight to hide memory latency; too few and the GPU stalls, too many and they starve each other of registers. Separately, a kernel is either memory-bound, limited by how fast it can read data, or compute-bound, limited by arithmetic throughput. Most large-language-model inference is memory-bound, which is why techniques that reduce data movement — fusion, quantization, better tiling — matter more than raw arithmetic. A good kernel writer is essentially a traffic engineer for bytes.

The practical takeaway

For a self-hoster, this explains a puzzle: two models with identical parameter counts can run at very different speeds, and the same model can speed up overnight when a serving engine ships better kernels. Choosing a runtime that invests in fast kernels for your specific hardware often matters more than shaving parameters.

Kernel launches themselves carry overhead, which is why serving engines work hard to reduce their number. Each time the CPU tells the GPU to run a kernel there is a fixed cost, and a model naively implemented as thousands of tiny separate launches can spend more time dispatching work than doing it. Techniques such as kernel fusion combine several logical operations into one launch, and CUDA graphs capture a whole sequence of launches so it can be replayed with almost no per-launch cost. For a self-hoster this is another reason the same weights can run at very different speeds across runtimes: a stack that fuses aggressively and minimizes launches keeps the GPU fed, while one that dispatches naively leaves it stalling between tiny jobs. It is also why the first tokens of a request can feel slower than the rest, as the machinery warms up.

Kernels are dispatched through CUDA and frequently target specialized Tensor Cores, and the matrix multiply most of them exist to accelerate is a GEMM under the hood.