What is a CUDA Kernel?

A kernel is the unit of CUDA code a programmer actually writes, the rough equivalent of a function in ordinary CPU programming. The difference is what happens when you call one: a kernel is launched once and returns once, but in between, it runs many times over, once per thread, with no guaranteed order between those runs and no guarantee they happen one after another rather than side by side on separate execution units.

The full set of threads running a single kernel launch is called a kernel grid, made up of thread blocks, the top level of the CUDA programming model's thread hierarchy. Because a grid spans however many Streaming Multiprocessors the launch needs, it operates at the scale of the whole GPU, and the memory it reaches at that scale is the device's global memory (what shows up on a spec sheet as GPU RAM).

In CUDA C++, a kernel function is handed raw pointers into that global memory when the host launches it, and it returns nothing. All of its work is done as a side effect: reading and writing through those pointers.

A worked example: multiplying two matrices

The standard first kernel anyone writes is square matrix multiplication, computing C = A x B. Two implementations below use the same textbook algorithm but map it onto the hardware differently, and the gap between them is most of what "writing a fast kernel" is about.

The naive version assigns one output element to one thread. Each thread walks a full row of A and a full column of B, multiplying and accumulating as it goes:

__global__ void matmul_naive(float* A, float* B, float* C, int N) {
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;

  if (row < N && col < N) {
    float acc = 0.0f;
    for (int k = 0; k < N; k++) {
      acc += A[row * N + k] * B[k * N + col];
    }
    C[row * N + col] = acc;
  }
}

Each thread does exactly one multiply and one add per pair of reads from global memory. That ratio is bad: the CUDA Cores' arithmetic throughput in FLOPs per second is far higher than the memory bandwidth between GPU RAM and the SMs, so this kernel spends most of its time waiting on loads rather than computing.

A tiled version fixes the ratio by staging square tiles of A and B into shared memory, so a thread block loads each value from global memory once and reuses it many times from the much faster on-chip memory:

#define TILE 16

__global__ void matmul_tiled(float* A, float* B, float* C, int N) {
  __shared__ float As[TILE][TILE];
  __shared__ float Bs[TILE][TILE];

  int row = blockIdx.y * TILE + threadIdx.y;
  int col = blockIdx.x * TILE + threadIdx.x;

  float acc = 0.0f;
  for (int t = 0; t < N / TILE; ++t) {
    As[threadIdx.y][threadIdx.x] = A[row * N + (t * TILE + threadIdx.x)];
    Bs[threadIdx.y][threadIdx.x] = B[(t * TILE + threadIdx.y) * N + col];
    __syncthreads();

    for (int k = 0; k < TILE; ++k) {
      acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
    }
    __syncthreads();
  }
  C[row * N + col] = acc;
}

Every pass through the outer loop loads two values per thread and then runs 16 multiply-adds against data already sitting in shared memory, a far better FLOP-to-load ratio than the naive version, for exactly the reason the roofline model predicts.

This still isn't a production-grade kernel. Simon Boehm's widely-cited "How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance" worklog walks through ten progressively faster versions of this same problem, and the two shown here sit near the start of that sequence. And both examples here only use CUDA Cores; the fastest matrix-multiply kernels on modern GPUs run on Tensor Cores instead, which trade flexibility for a much higher arithmetic ceiling on exactly this kind of operation.

You don't have to hand-write a tiled kernel to get tiled-kernel performance. Vendor libraries like cuBLAS and CUTLASS already implement heavily-tuned matmul kernels for every recent GPU generation, and Aquanode's marketplace gives you H100, H200, B200 and MI300X instances to run them on without provisioning your own CUDA toolchain first.