Two kernels launched back to back in the same stream run strictly in order: the second does not begin until the first has retired, even though much of the second kernel’s early work does not depend on the first kernel’s output. Loading launch parameters, initializing shared memory, and computing addresses – the kernel’s preamble – could in principle run while the first kernel is still draining. For a long chain of small kernels, such as the layers of an inference network or the steps of an iterative solver, that strict serialization exposes each kernel’s launch latency and preamble after the previous kernel’s tail, one kernel at a time.
Programmatic dependent launch (PDL), added with Hopper (SM
9.0), lets the dependent kernel begin its preamble while the primary
kernel is still finishing. The primary kernel calls
cudaTriggerProgrammaticLaunchCompletion() once it has produced enough
state for the dependent kernel to start; the dependent kernel runs its
preamble, then calls cudaGridDependencySynchronize() at the point where
it first reads the primary’s results, blocking there until the primary
has fully retired. The dependent kernel is launched with
cudaLaunchKernelEx() and a launch attribute that enables programmatic
serialization, so the runtime overlaps the two grids rather than
serializing them. Figure 7-6 places both calls on the timeline of the
two overlapping grids.
Figure 7-6. Timeline of a programmatic dependent launch.
The tail of one grid then overlaps the preamble of the next. The gain is lower latency rather than higher throughput, and it grows as the kernels shrink and the chain lengthens – the regime where launch and preamble overhead, not the kernels’ own work, sets the pace.
The trigger is optional. If the primary kernel never calls
cudaTriggerProgrammaticLaunchCompletion(), the completion signal fires
on its own as the primary’s last warps retire; the dependent then starts
only as the primary finishes, forfeiting the overlap but computing the
same result. Correctness rests instead on the dependent’s
cudaGridDependencySynchronize(), which blocks until the primary has
fully retired and flushed its results to global memory. Because the
dependent’s preamble runs alongside the primary’s tail, the
synchronization must precede the dependent’s first read of the primary’s
output; a read placed ahead of it races against data the primary has not
yet produced. The dependency that ties the two grids together is
declared on the host, in the dependent’s launch, so the primary and
dependent are designed as a producer and consumer pair rather than
composed after the fact.