Kernels, grids and threads

CUDA is how you turn the GPU's SMs and warps into a program. You will not rewrite vLLM in raw CUDA tomorrow, but every serving mystery — why a tiny elementwise kernel kills decode, why a matmul needs tiling — starts with the launch model.

The big idea

One function, many workers. You write the body for one index. The hardware runs it across a grid of blocks of threads. A block is scheduled onto one SM. Threads in a block share shared memory and can __syncthreads().

A restaurant kitchen with one recipe and 256 line cooks is the picture. It breaks at memory: the cooks do not share a fridge with the dining room. Data must be copied (or mapped carefully) between host and device.

grid  → many blocks
block → many threads (grouped into warps of 32)
thread → one lane with its own registers and threadIdx

Walking through vector add

Each thread computes one index. This is the shape of almost every data-parallel kernel.

__global__ void vec_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
int threads = 256;
int blocks = (n + threads - 1) / threads;
vec_add<<<blocks, threads>>>(d_a, d_b, d_c, n);

For n = 1_000_000 and 256 threads/block:

blocks = ceil(1e6 / 256) = 3907
total threads launched ≈ 3907 × 256 = 1,000,192

Extra threads in the last block hit the i < n guard and idle. That idle is normal. Under-launching (too few threads for the GPU) is worse.

The if (i < n) at the end of a grid is mostly fine: only the last warp partially predicates. Data-dependent branches inside a hot loop — different expert paths, per-thread early exit — serialize warp execution. Structure algorithms so warps stay uniform.

One practical check: print blockDim, gridDim and the problem size. Off-by-one grids that leave most of the GPU idle are common when porting a tutorial sized for n=1024 to production tensors.

Common mistakes

• Forgetting the i < n guard. The last block overruns the array.
• Launching too few threads. The SMs sit idle while you think the kernel is "running."
• Assuming host and device share memory. They do not, unless you opted into unified memory — and serving should not.

Check yourself

1. n = 1,000,000, 256 threads/block. How many blocks? Answer: ceil(1e6 / 256) = 3907.
2. Why is a little idle in the last block OK? Answer: those threads hit i < n and wait. Under-launching the whole grid is worse.

Key takeaways

• A kernel is a device function. The launch maps a grid of blocks of threads onto SMs.
• Index with blockIdx * blockDim + threadIdx, and guard the tail.
• Host and device memories are separate. Global memory is HBM, the slow path.

Further learning

Next, make those loads cheap in coalescing and tiling.