What a Kernel Launch Physically Does

Launching a kernel looks like calling a function, and that resemblance causes most of the early confusion. Nothing is called. A description of the work is written into a queue, your CPU carries on immediately, and at some later moment the GPU picks that description up and starts handing pieces of it to hardware. Understanding those two facts, that it is a queue and that it is asynchronous, explains a surprising number of bugs.

When you launch a kernel, the runtime packages the grid dimensions, the block dimensions, the kernel's address and its arguments into a command, and appends that command to a stream. A stream is an ordered queue: commands in the same stream run one after another, commands in different streams may overlap. The call then returns. It does not wait for the kernel to start, let alone finish.

On AMD hardware that command is not an abstraction; it is a 64-byte record with a published layout, the AQL kernel dispatch packet from the HSA specification. The runtime copies your argument values into a separate buffer, then writes the packet into a ring buffer in memory the GPU can read, and finally "rings the doorbell": a single store to a special address that tells the GPU new work is waiting. No system call, no driver round trip, which is why a launch is cheap at all.

Nvidia's equivalent is not publicly documented to the same level, but the shape is the same: a command in a memory-resident queue, and a doorbell.

On the device, a command processor reads from the queues, and a piece of hardware usually called the work distributor begins handing thread blocks out to SMs (Nvidia) or CUs (AMD) that have room for them. Your grid is simply a count of blocks to be placed. Nothing about the launch guarantees when any of them runs.

This page is about that mechanism. For the syntax itself, the triple chevrons, hipLaunchKernelGGL and the index arithmetic inside the kernel, see CUDA & HIP.

Because the launch returns immediately, the host can run ahead. That is the point: you want the CPU queueing more work, or preparing the next batch, while the GPU is busy. It also means three things beginners hit almost immediately.

The fix for all three is the same: know where your synchronization points are. For timing, record events into the same stream on either side of the kernel and read the elapsed time between them, or let a profiler do it. The code sample further down shows the difference on real hardware.

The work distributor assigns each thread block to exactly one SM/CU, and it stays there for its whole life. It is never split across two, and it never migrates. That single rule explains several things that otherwise look arbitrary.

Because a block lives entirely on one SM/CU, the threads in it can share that SM/CU's on-chip scratchpad and can synchronize with a barrier. Both facilities exist only within a block, because only within a block is the hardware guaranteed to be the same. It also means a block must fit: its registers, its scratchpad allocation and its thread count all have to be available on one SM/CU, and if you ask for more than one SM/CU can provide, the launch fails rather than being split.

Once it lands, the block is cut into fixed-size groups of consecutive threads: warps of 32 on Nvidia, wavefronts of 64 on AMD's datacenter GPUs. That group, not the thread and not the block, is what the hardware actually schedules: one instruction is issued for all of its threads at once. A block of 256 threads is therefore 8 warps or 4 wavefronts. A block of 100 threads still occupies 2 whole wavefronts, and 28 of those 128 lanes do nothing for the entire kernel, which is why block sizes are chosen as multiples of the group size.

The corollary is that the number of resident blocks per SM/CU is a resource calculation, not a scheduling preference. That calculation is what occupancy measures.

Put numbers on it. The card is an AMD MI300X with 304 CUs. Each CU has 4 SIMD units, and each SIMD can keep up to 8 wavefronts resident, so a CU holds at most 32 wavefronts, or 2,048 threads, at a time. The kernel processes 1,048,576 elements, one per thread, in blocks of 256 threads.

"Light kernel" matters. Each SIMD has 512 vector registers to share among its resident wavefronts, so 8 wavefronts fit only if each thread uses 64 or fewer. A heavier kernel using 128 registers fits 4 wavefronts per SIMD, 16 per CU, so only 4 blocks per CU and 1,216 across the GPU. The same 4,096 blocks now take 4 waves, and the last one is only 37% full: for that final stretch, most of the GPU sits idle waiting for a few hundred blocks to finish. This effect is called wave quantization, or the tail effect, and it is why a grid slightly larger than a whole number of waves can be noticeably slower than one slightly smaller.

The opposite mistake is too few blocks. Launch 256 blocks on this card and 48 of its 304 CUs get nothing at all, while the other 256 each run a single block, 4 of their 32 wavefront slots, with little else to switch to when those stall on memory. As a rule, a grid should have several times more blocks than the GPU has SMs/CUs.

There is no ordering guarantee between blocks. Block 7 may finish before block 0 starts. Blocks may run concurrently, or one after another, depending entirely on how many fit and how many SMs/CUs are free. A grid larger than the machine simply runs in waves as earlier blocks retire.

So a correct kernel cannot assume any relationship between blocks, and there is no barrier that waits for all of them. If you need one, you end the kernel: a kernel boundary is the grid-wide synchronization point. This is also why reductions are written as a tree across multiple launches, or with atomics, rather than as one pass with a global barrier.

Within a block, ordering is available and cheap. Across blocks, it is not available at all. Designing around that split is most of what makes GPU algorithms look unusual.

Enqueueing a kernel costs on the order of a few microseconds of CPU time, and there is further latency before the work actually begins on the device. That is irrelevant for a kernel that runs for milliseconds and ruinous for one that runs for two microseconds.

The 5 µs is a round number; the sample below measures it on your own card. Launches in a stream can overlap with the previous kernel still running, so a long queue hides some of this, but a CPU that cannot enqueue faster than the GPU drains will leave the GPU idle between kernels.

This is why fusing several tiny kernels into one larger kernel is such a reliable optimization, and why frameworks that issue thousands of small operations per step spend a startling fraction of their time simply launching them. When the shape of the work is fixed and repeated, both vendors offer a graph API that records a whole sequence of launches once and replays it with a single submission.

The companion sample runs three small experiments in one HIP program: launch cost measured with 10,000 empty kernels, a kernel timed wrongly with the host clock and correctly with events, and a kernel whose blocks each take a ticket from an atomic counter as they start, to show the order they really ran in. The heart of the timing experiment:

auto t1 = Clock::now();
hipEventRecord(start);
busyKernel<<<grid, block>>>(d_data, n, iters);
hipEventRecord(stop);
double wrongUs = microsSince(t1);   // no sync yet: this timed the enqueue
hipEventSynchronize(stop);
hipEventElapsedTime(&rightMs, start, stop);   // this timed the kernel
            

On a small Nvidia card, built as CUDA, it printed the following. An empty launch cost about 2 µs. The host clock reported 39 µs for a kernel the events measured at 19.9 ms, a factor of 500. And of 4,096 blocks, 3,648 did not start in index order; the first few started as 3 4 5 0 1 2 9 10 11 6 7 8. Your numbers will differ; the shape will not.

Profilers make the gap between "launched" and "running" visible, which is the fastest way to believe it.

• Make block sizes a multiple of 64. That is a whole number of wavefronts on AMD and of warps on Nvidia; 256 is a sensible default.
• Launch several times more blocks than the GPU has SMs/CUs, so every one of them has work and something to switch to.
• Never assume anything about the order blocks run in. If blocks must wait for each other, end the kernel and start another.
• Time with events or a profiler, never with a host clock around the launch alone.
• Check the launch's error code, but remember that faults inside the kernel appear only at the next synchronization.
• If a kernel runs for only a few microseconds, suspect launch overhead before suspecting the kernel.