CUDA Kernel Execution: From Grid Launch to SM Execution Units

Hello! This course begins by building a performance model from the hardware upward. Later lessons will use that model to explain measurements, optimize kernels, and reason about profiling infrastructure. Here, the task is narrower: trace one kernel launch from the grid requested by host code to the instructions issued on an SM.

The distinction to keep in mind is between logical work and physical execution. A launch specifies threads and blocks; it does not specify which SM receives each block, when that block runs, or which execution unit handles each instruction.

A launch defines a grid, not a hardware placement

Consider this CUDA C++ kernel:

__global__ void 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];
    }
}

// Host code, assuming a, b, and c refer to suitable device-accessible data:
int threads_per_block = 128;
int blocks = (n + threads_per_block - 1) / threads_per_block;
add<<<blocks, threads_per_block>>>(a, b, c, n);

The values inside <<<...>>> are the execution configuration: first the number of blocks in the grid, then the number of threads in each block. Suppose $n=1000$. The host requests eight blocks of 128 threads, or 1,024 logical CUDA threads. It has not requested 1,024 simultaneously executing operations.

Watch how a launch and its thread IDs are introduced in the following lecture excerpt. The lecturer occasionally compares CUDA with another programming system; that comparison is not needed here.

Inside add, every thread evaluates the same index expression using its own blockIdx.x and threadIdx.x. For example, thread 45 of block 6 computes $6\cdot 128+45=813$, so it handles element 813. The final block contains threads whose indices extend beyond 999. The if guard prevents those threads from accessing outside the arrays.

The guide makes the built-in identifiers precise and explains how multidimensional thread indices are linearized:

Grids and blocks can each have one, two, or three dimensions. Those dimensions help express the shape of the work, but they do not turn grid coordinates into SM numbers. The index calculation assigns data ownership to a thread, not hardware placement.

Blocks are placed on SMs

A GPU contains multiple streaming multiprocessors (SMs). For this lesson, think of an SM as the hardware location that can hold a block’s execution state and issue its instructions. SMs contain registers and functional units, among other resources; the details vary by architecture.

The following diagram distinguishes the grid of requested blocks from the subset currently active on SMs.

A block executes on one SM, rather than having some of its threads sent to one SM and the rest to another. An SM may have multiple blocks active at once if it has enough resources for them. If no SM can admit another block, remaining blocks wait until capacity becomes available. Crucially, active means that a block’s state is resident; it does not mean that all its threads perform an operation in every cycle.

For the eight-block add grid, you therefore cannot infer that block 6 runs on SM 6, or even that block 6 starts after block 5. The kernel is written so that its blocks can run in any order: each computes its own output elements. This independence is what lets the same grid scale across devices with different numbers of SMs.

There are two scheduling decisions to keep separate. Block placement determines which blocks become resident on which SMs. Once a block is resident, warp scheduling within that SM determines which of its ready instructions get issued. The next section identifies those warps.

A resident block consists of warps

Hardware organizes a block’s threads into warps of 32. With the one-dimensional, 128-thread block above, each block has four warps: local thread indices 0–31, 32–63, 64–95, and 96–127. Within a warp, a thread has a lane numbered 0–31. Thus, thread 45 of block 6 is in that block’s second warp and has lane 13.

Read the guide’s description of what it means for a warp to execute a kernel:

SIMT—single-instruction, multiple-threads—lets you write a kernel as though each CUDA thread has its own indices, values, and control flow. At issue time, a warp groups work on an instruction across its participating lanes. If threads take different paths, not every lane participates in every instruction. Treat “the warp executes together” as a useful programming-model description, not a claim that every thread always has the same control-flow state or that an SM necessarily has 32 physical units for every instruction type.

Return to the last block of add. Its fourth warp covers global indices 992–1023. Only indices 992–999 satisfy i < n: eight lanes perform the guarded loads, addition, and store; the other 24 do not perform those operations. All four warps of that block still need execution state on its SM. Detailed performance effects of divergence come later; for now, the important point is that a warp is the grouping used to issue work, while an individual lane determines whether it participates.

From a warp instruction to an execution unit

Once block 6 is resident, its four warps are available to scheduling hardware on that SM. A warp scheduler selects an instruction that a warp is ready to issue; the instruction is dispatched to hardware appropriate for its operation. For add, the work includes integer index or address calculations, memory loads, floating-point addition, and a memory store. These are not all performed by one undifferentiated “CUDA core.”

The lecture excerpt below uses an architecture-specific illustration to show execution contexts and several kinds of units. Its counts and lane widths describe the illustrated design, not a rule to apply to every GPU generation.

A useful distinction for interpreting a trace is resident versus issuing. An SM can retain state for many warps while issuing instructions from only a subset at a given moment. A warp’s source values may live in registers, but an arithmetic instruction still has to be issued to a suitable pipeline; a load or store uses memory-access hardware. The order and precise instruction sequence produced from C++ can differ from the simple source-level description.

One complete trace

For $n=1000$, follow the thread we identified earlier without assuming a particular SM number:

1. Host launch: add<<<8, 128>>> requests one grid containing eight 128-thread blocks.
2. Block placement: the device assigns block 6 to an SM when that SM can admit it. Other blocks may be active there or elsewhere; their order is unspecified.
3. Warp formation: block 6’s 128 threads form four warps. Local thread 45 belongs to the warp containing threads 32–63 and occupies lane 13.
4. Instruction issue: scheduling hardware on the block’s SM issues instructions for that warp as they become ready. Lane 13 computes global index 813, accesses a[813] and b[813], adds their values, and writes c[813] through the relevant arithmetic and memory pipelines.
5. Completion: the block finishes after its threads have completed their work. Its SM resources can then be used by further blocks, if any remain.

This trace moves between levels of description: the grid is the complete request; the block is the unit placed on an SM; the warp is the group used for instruction issue; and the execution units perform particular operations. None of these levels should be mistaken for a fixed one-to-one mapping to the next.

Takeaways

A CUDA launch specifies logical dimensions, while the device decides block placement. Each block stays on one SM, its threads are organized into 32-thread warps, and an SM issues warp instructions to operation-specific hardware. Residency is not the same as simultaneous execution, and neither block coordinates nor thread indices reveal physical placement.

In the next lesson, we will examine why an SM keeps multiple warps available: when one cannot issue its next instruction yet, scheduling another can help keep execution units busy.