GPU Kernels - The CUDA Execution Model

One of the hardest parts for me, as someone trying to learn about CUDA, GPU kernels and all things in between, is building real intuition of how everything fits together. Grids, blocks, threads and warps as they appear in the code, versus SMs, resident blocks, as they exist on the hardware and how the two actually map onto each other. In this concise post, I’ll go over these “elements” in the way that helped me the most understand them.

Two views that need to line up

• Program view (what you write / launch): grid → blocks → threads. E.g. kernel<<<num_blocks, threads_per_block, shared_mem_bytes>>> launches a grid of num_blocks blocks, each with threads_per_block threads. This is the abstraction the code is written against — blockIdx.x, threadIdx.x, blockDim.x are all program-view coordinates, and nothing about them mentions physical hardware.
• Hardware view (what actually executes it): GPU → SMs (streaming multiprocessors) → resident blocks → warps. Blocks get scheduled onto SMs; an SM can hold multiple blocks resident at once, up to whatever its register/shared-memory budget allows; each resident block’s threads get split into warps of 32 for actual execution.

Individual elements worth going deeper

• A block is a program-view unit of cooperation (shared memory, __syncthreads()) — but it only makes sense because all its threads land on the same SM. Threads in different blocks can’t cooperate this way even if they’re logically “next to each other” in the grid.
• Warps are a hardware-view execution unit, invisible in the code but central to performance — directly related to “coalescing”.
• Occupancy is the hardware-view constraint that decides how many blocks (and therefore how many warps) can be resident on one SM simultaneously — this is where threads_per_block in the kernels feeds back into a hardware limit rather than just a program-view launch parameter.

Visualizing everything together

A simple, introductory way of visualizing everything working together is described by the following sketch: