What Is a CUDA Kernel? A Visual Explainer

The thing your favorite LLM is doing right now, billions of times per second, is running CUDA kernels. They are the smallest unit of work the GPU understands. They are also the place where every interesting AI system in 2026 lives or dies.

This is a visual explainer. We will start from the single-most-important idea in GPU programming, which is the difference between thinking serially and thinking in parallel. Then we will build up the kernel: what it is, what its launch looks like, how threads find their work, where the speed comes from, and what trips it up. By the end you should be able to read a kernel and have an opinion about whether it is fast.

The Parallel Revelation

A CPU is a small number of strong workers. Eight cores, sixteen threads, each running at 4 GHz, each capable of complex branching and big caches. A GPU is the opposite. An NVIDIA H100 has 132 streaming multiprocessors (SMs), each able to have 2,048 threads in flight. That is over a quarter of a million threads, each running at a slower clock, each individually less capable than a CPU thread. The trade is intentional. If you have a task that decomposes into independent pieces, the GPU will eat it.

Almost everything in deep learning decomposes that way. A matrix multiply is independent dot products. A convolution is independent windows. Attention is independent rows of a softmax. All of these are jobs that look like "do the same thing, everywhere, at once." That is exactly what a kernel expresses.

A Kernel, In One Sentence

A CUDA kernel is a function that runs on the GPU and is executed in parallel by many threads.

That sentence has the entire model. The function is just C++ (with a few extra keywords). The runtime is "in parallel by many threads". You write the function once. The GPU runs it as many copies as you ask for, all at the same time, each on slightly different data.

The function is marked with __global__, which tells the compiler "this gets compiled to GPU instructions and is launched from the host". To launch it, you use the triple-chevron syntax that has confused every CUDA beginner since 2007.

The Anatomy of a Launch

The chevrons hold the execution configuration: how many blocks, how many threads per block. Everything inside the parentheses is the regular argument list. There is nothing else to it. The compiler turns the chevron call into a runtime API call, the runtime ships your function and its arguments across the PCIe bus to the GPU, and the GPU schedules the work.

Two things are doing all the work in that line. The first is blockDim: how many threads make up a block. A block is the unit that is guaranteed to run together on a single SM and can talk to each other through shared memory. The second is gridDim: how many of those blocks you want. The grid is the entire population of threads for the launch. Multiply them and you get the total thread count. The standard ceiling-division trick to cover an array of size N is (N + blockSize - 1) / blockSize blocks of blockSize threads each.

Threads, Blocks, Grids

Three levels of hierarchy. Thread is the smallest unit, the actual instance of your function. Threads in the same block live on the same SM and can synchronize and share memory cheaply. Blocks make up the grid, and they cannot directly talk to each other (without going through global memory or, on Hopper, distributed shared memory inside a block cluster).

Inside the kernel, every thread can ask "who am I?" and gets back four built-in variables: threadIdx, blockIdx, blockDim, and gridDim. The standard way to convert that identity into an index into a global array is the most important formula in the language.

That formula is doing the work of the for loop in a CPU implementation. There is no for loop in a kernel. The for loop is the launch itself. Each thread is one iteration, running concurrently with all the others. This shift, from "one worker iterating" to "every iteration is its own worker", is the change of frame that makes everything else click.

The Memory Hierarchy

The threads themselves are cheap. Feeding them is the entire problem.

A modern GPU is structured as a deep memory hierarchy. The closer to the SM, the faster and the smaller. The further out, the larger and the slower. Writing a kernel is, more than anything else, an exercise in moving the right data into the right tier and reusing it as much as possible before sending it back out.

A useful mental model: HBM is the warehouse, L2 is the loading dock, shared memory is the workbench, registers are your hands. You do not assemble things in the warehouse. You bring a tile to the workbench, you finish your task there, and you ship the result back. Most kernel optimization is some version of "keep more of the work on the workbench, make fewer trips to the warehouse".

The numbers above are approximate H100 figures. Blackwell B200 raises HBM to roughly 8 TB/s and adds FP4 Tensor Cores, but the shape of the hierarchy is the same. The ratios that matter (registers ~10× faster than shared, shared ~3× faster than HBM) hold across generations.

Warps and the Cost of Branches

One more piece of the architecture you need before you can read a real kernel. Threads are not actually scheduled individually. The SM groups them into bundles of 32 threads called warps, and the warp is the unit it executes. Every cycle, every thread in a warp issues the same instruction at the same time, on its own data. NVIDIA calls this SIMT: Single Instruction, Multiple Threads.

The neat consequence: launching threads in multiples of 32 is basically free. Below 32, you are paying for unused lanes in the warp. The annoying consequence is what happens at branches.

When the threads in a warp disagree about which way to go in an if, the warp serializes. It runs the true branch with the false-branch threads masked off, then runs the false branch with the true-branch threads masked off. Both paths cost time. This is warp divergence, and it is the most common reason a kernel that "should be fast" turns out to be merely fine.

The cure is structural: arrange your data and your branches so that whole warps go the same way. Sort your work into groups before launching. Pad to the warp boundary. Use predicated math (compute both sides and select) instead of branches when the true and false paths are short. The fastest kernels in the world have almost no branches inside their inner loops.

A Complete Kernel: SAXPY

Time to put it together. SAXPY stands for single-precision A times X Plus Y, and it has been the "hello world" of GPU programming since the 1980s. The operation is y[i] = a * x[i] + y[i] for every i. It is embarrassingly parallel. Every thread does one multiplication and one addition, with no dependencies on any other thread.

The kernel is six lines. The launch is one line. That is the entire program. Every thread reads two floats from HBM, does one fused multiply-add, writes one float back. There are no loops inside the kernel. The for loop is the grid.

SAXPY is also a useful sanity check for hardware. On an H100, this kernel runs at about 3 TB/s, almost exactly the HBM3 bandwidth, because every thread reads 8 bytes, writes 4 bytes, and does 2 floating-point operations. It is memory-bound. You cannot make it faster by adding more compute. You can only make it faster by moving less data, which is exactly the move that FlashAttention pulled off for the attention operation.

Why Every LLM Lives on Kernels

Once you can write SAXPY, you have the model for everything else. The trick is that the operations that matter are harder. A matrix multiply has data dependencies. Softmax has a reduction. Attention has both, plus a giant intermediate matrix you do not want to materialize. These are all kernels too, just much more carefully designed ones.

Three pieces of work define the modern landscape.

FlashAttention

FlashAttention (Tri Dao, 2022, with v2 in 2023 and v3 in 2024) is the single most important AI kernel of the last five years. The insight: the attention output for one query block can be computed by streaming through the keys and values in tiles, keeping a running softmax in shared memory, and never writing the full attention matrix to HBM. Memory traffic drops from O(N²) to O(N). On long contexts, the speedup is several×.

FlashAttention-3 added warp-specialization (different warps in a block do different jobs, like loading vs computing), asynchronous WGMMA on Hopper Tensor Cores, and FP8 support. It hits 740 TFLOPs/s in FP16 and 1.2 PFLOPs/s in FP8, around 75% of an H100's theoretical peak. Without this kernel, the 128K and 1M context windows that GPT-4 and Llama 3 ship with would not be commercially viable.

Triton

Triton (Phil Tillet, OpenAI, 2019) is a Python-like language that compiles to PTX. It handles a huge amount of the bookkeeping (tile loading, register allocation, pipelining) for you. It does not always match a hand-written CUDA kernel on the latest hardware, but it is fast enough that most of PyTorch's newer fused ops ship as Triton kernels. If you want to write a custom kernel without spending three weeks learning CUTLASS, Triton is the right entry point.

ThunderKittens

ThunderKittens (Hazy Research, Stanford, 2024) is the cleanest abstraction in the space. Its unit of work is a 16×16 tile (the size that fits the H100 Tensor Core), and its primitives map directly to warp-level instructions. Their FlashAttention implementation is under 100 lines of code and reaches 855 TFLOPs/s, 86% of H100 peak, matching FA3. They have since produced kernels for Mamba-2, linear attention, RoPE, LayerNorm, and FlashFFTConv, several of which beat the best Triton implementations by 6-14×.

The lesson across all three: the abstraction that wins is the one that exposes the GPU's actual structure (warps, tiles, async memory copies) instead of hiding it. PyTorch hides too much for frontier kernels. Pure CUDA shows too much. Triton, ThunderKittens, and CUTLASS each find a different equilibrium in the middle.

How to Read Kernel Performance

When you run a kernel under Nsight Compute, the first question to ask is whether you are memory-bound or compute-bound. The roofline model gives you a clean answer: plot achieved performance against arithmetic intensity (FLOPs per byte loaded). If you are below the diagonal "memory roof", you are waiting on HBM. If you are below the horizontal "peak FLOPs roof", you are waiting on the Tensor Cores.

SAXPY has an arithmetic intensity of 0.17 FLOPs/byte, far below where compute could ever bottleneck it. It is unfixably memory-bound. A matmul with a big M, N, K has arithmetic intensity in the hundreds, comfortably in the compute-bound regime, and that is why GEMMs are how we measure peak FLOPs in the first place.

The other number that matters is occupancy: how many of the SM's 2,048 thread slots are actually in flight. Low occupancy is usually a sign of register pressure (each thread using too many registers) or shared memory pressure. The fix is almost always to use smaller tiles, fewer registers, or both.

That is the entire mental model: arithmetic intensity tells you which roof you are hitting, occupancy tells you whether you are even on the roof. Everything else (warp efficiency, memory coalescing, shared memory bank conflicts) is a way to climb a little higher.

The interesting work in 2026 is happening just below those rooflines, where 75% of peak is the new floor for serious attention kernels and 90% is the active research frontier. The tools to write those kernels are better than ever. The hardware (Hopper, Blackwell) is finally fast enough that getting the kernel right matters more than getting more silicon. If you build AI infrastructure and you have not stared at an ncu profile in a while, this is the year to start.