Markdown

Dynamic parallelism and device-initiated launch¶

Scope: launching work from the device, via CUDA Dynamic Parallelism (a kernel launches kernels) and device graph launch (a kernel launches a preinstantiated CUDA graph), to keep orchestration on the GPU and cut CPU launch latency for data-dependent work such as MoE routing and variable-sized tasks. Also covers Programmatic Dependent Launch (PDL), a related but distinct host-launched overlap mechanism that lets a second kernel's non-dependent work start before the first kernel it follows has fully completed.

What it is¶

Normally every kernel and every graph is launched by the host: the CPU decides what runs next, configures the launch, and submits it across the PCIe/NVLink boundary. For data-dependent control flow (where the next launch's shape or even existence depends on values the GPU just computed), this forces a device-to-host round trip per decision. The GPU stalls waiting for the CPU to read a result and dispatch the follow-up.

Two device-side mechanisms move that decision onto the GPU:

CUDA Dynamic Parallelism (CDP) lets a running kernel launch child kernels using the same triple-chevron syntax as host code: kernel<<<Dg, Db, Ns, S>>>(args), where Dg/Db are grid/block dims, Ns is dynamic shared memory, and S is a stream. Each thread can make an independent launch decision; device-side launches are asynchronous with respect to the launching thread, exactly like host launches. (NVIDIA CUDA Programming Guide, Dynamic Parallelism)

Device graph launch lets a running kernel launch an entire preinstantiated CUDA graph (a fixed DAG of kernel/memcpy/memset/child-graph nodes) based on runtime data. The graph is built and instantiated once on the host, uploaded to the device, then relaunched from device code with near-flat latency regardless of graph width. (NVIDIA Technical Blog, Enabling Dynamic Control Flow in CUDA Graphs with Device Graph Launch)

Both keep the orchestration loop GPU-resident. The CPU is removed from the per-decision critical path. (Fregly, Ch. 12)

A third, easily confused mechanism, Programmatic Dependent Launch (PDL), is not device-initiated launch at all: both kernels are still queued by the host, in the same stream, in program order, exactly as with any ordinary back-to-back launch. What PDL changes is only whether the second kernel's thread blocks are allowed to start executing before the first kernel's grid has fully finished, so the second kernel's non-dependent prologue work (e.g. zeroing buffers, loading constants) overlaps the first kernel's tail instead of waiting behind it. See "Programmatic Dependent Launch" below for the full mechanism; it is covered on this page because it answers the same "how do I overlap two dependent kernels" question CDP and device graph launch answer, via host launch order rather than a device-initiated one.

Why use it¶

The cost being eliminated is launch latency on the device-to-host-to-device path, not compute. When a kernel must inspect its own output to decide the next launch, host orchestration pays a full round trip (submit, GPU completes, CPU observes, CPU dispatches) for every decision. On launch-bound, fine-grained, data-dependent workloads this dominates wall-clock time.

Device graph launch is the stronger of the two for repeated structured work: its launch latency is lower than host launch and stays roughly flat as the graph grows wider, where host launch cost scales with the number of nodes submitted. (NVIDIA Technical Blog, Device Graph Launch) The book reports roughly 2x lower launch latency versus host-side launch for the same graph; treat that as illustrative, not a target for your hardware. (Fregly, Ch. 12)

Concrete motivating cases:

MoE routing: a router kernel computes per-token expert assignments, then the GPU itself must dispatch the right expert kernels with token-count-dependent shapes. Keeping the dispatch on-device avoids a host round trip per routing step.
Variable / recursive work: tree and graph traversal, adaptive mesh refinement, sparse expansion, where the amount of follow-on work is unknown until the parent runs.
GPU-resident scheduler loops: a graph that relaunches itself (tail launch) to drive an iterative pipeline without returning to the CPU between iterations. (Fregly, Ch. 12)

When to use it (and when not)¶

Reach for device-initiated launch when all of the following hold:

The next launch's parameters or existence is data-dependent on values the GPU just produced.
Profiling shows the host is on the critical path, with visible device-to-host-to-device gaps per decision (confirm with a Nsight profiling pass).
The decision granularity is fine: many small dependent launches, so the round-trip cost is paid often.

Prefer device graph launch over raw CDP when the dependent work is a fixed-structure DAG replayed with varying data: you get the graph's batched scheduling plus device launch, and the node set is validated once at instantiate.

Do not use these when:

The pipeline is static and host-known. Plain CUDA graphs with host replay are simpler and sufficient.
Kernels are large and already saturate the GPU. Launch latency is noise; there is nothing to amortize. Check the roofline / arithmetic-intensity regime first.
The work can be expressed as a single grid with persistent kernels or grid-stride loops. A persistent kernel that pulls work from a queue often beats per-item child launches and avoids device-launch overhead entirely. (Fregly, Ch. 12)
You need the parent to block on and read child results mid-kernel. As of CUDA 12.0 that pattern is gone (see below); restructure as a tail launch instead.

CDP carries real per-launch overhead (each device launch allocates from a device-side launch pool and consumes a pending-launch slot). Over-decomposing into millions of tiny child grids regresses performance. Launch coarse-grained children, or use a persistent kernel. (Fregly, Ch. 12)

Architecture¶

flowchart TD
  H["Host builds and uploads graph"] --> S["Scheduler kernel"]
  S --> D{"Runtime decision on GPU"}
  D -->|"irregular child grid"| C["CUDA Dynamic Parallelism"]
  D -->|"fixed uploaded DAG"| G["Device graph launch"]
  C --> T["Tail-launch consumer"]
  G --> R["Tail relaunch or fire-and-forget"]
  P["Host queues PDL primary and secondary"] --> O["Opportunistic prologue overlap"]

CDP creates a child grid at runtime. Device graph launch replays an executable graph instantiated and uploaded by the host. PDL is separate: the host launches both kernels and only their permitted execution overlap changes.

How to use it¶

Build flags (CDP and device graph launch)¶

Device-side launch requires relocatable device code and the device runtime library:

nvcc -arch=sm_90 -rdc=true device_launch.cu -o device_launch -lcudadevrt

-rdc=true: relocatable device code (required to call the device runtime).
-lcudadevrt: link the CUDA device runtime.

(NVIDIA CUDA Programming Guide, Dynamic Parallelism)

CDP2: no mid-kernel device-side synchronize¶

The legacy cudaDeviceSynchronize() inside device code (CDP1), which let a parent block until its children finished, was deprecated in CUDA 11.6 and removed in CUDA 12.0 (CDP2). On compute capability 9.0+ only CDP2 exists; CDP1 is not available. For compute capability < 9.0 you can opt back into the old behavior at compile time with -DCUDA_FORCE_CDP1_IF_SUPPORTED, but do not build new code on it. (Host-side cudaDeviceSynchronize() is unaffected and still supported.) (NVIDIA CUDA Programming Guide, Dynamic Parallelism; NVIDIA Developer Forums, cudaDeviceSynchronize from device code is deprecated)

The CDP2 replacement pattern: instead of blocking on a child and consuming its result inline, launch the consumer as a tail-launch child into cudaStreamTailLaunch. Tail launches execute only after the launching grid and its fire-and-forget children complete, so the consumer sees the producer's writes without an in-kernel sync. (NVIDIA CUDA Programming Guide, Dynamic Parallelism)

Device-code named streams for CDP:

cudaStreamFireAndForget: child runs independently; parent does not wait.
cudaStreamTailLaunch: child runs after the parent grid (and its fire-and-forget children) complete.

How to develop with it¶

CDP example: data-dependent child launch (CDP2)¶

#include <cuda_runtime.h>

__global__ void process_expert(const float* tokens, int count, int expert_id);

__global__ void reduce_results(const float* partials, int n_experts);

// Router decides, per block, how much follow-on work to launch.
__global__ void route(const float* tokens, const int* expert_counts,
                      const int* expert_offsets, int n_experts,
                      float* partials) {
    if (threadIdx.x == 0 && blockIdx.x == 0) {
        for (int e = 0; e < n_experts; ++e) {
            int count = expert_counts[e];           // computed earlier on-device
            if (count == 0) continue;               // skip empty experts
            int blocks = (count + 255) / 256;
            // Fire-and-forget: independent expert kernels, sized at runtime.
            process_expert<<<blocks, 256, 0, cudaStreamFireAndForget>>>(
                tokens + expert_offsets[e], count, e);
        }
        // Tail launch: runs only after all fire-and-forget children finish,
        // so it observes their writes without an in-kernel synchronize (CDP2).
        reduce_results<<<1, 256, 0, cudaStreamTailLaunch>>>(partials, n_experts);
    }
}

Each process_expert child is sized from a count the GPU computed, with no host round trip. The reduce_results tail launch replaces the removed cudaDeviceSynchronize() + inline reduce. (NVIDIA CUDA Programming Guide, Dynamic Parallelism)

Memory consistency: parent and child share global memory with weak ordering. A child's global writes are guaranteed visible to the parent only at a tail-launch boundary; local and shared memory are private and never shared across the parent/child boundary. Pass data through global memory only. (NVIDIA CUDA Programming Guide, Dynamic Parallelism)

Device graph launch: instantiate, upload, relaunch from device¶

A graph launched from the device must be:

Instantiated with the device-launch flag: cudaGraphInstantiate(&exec, graph, cudaGraphInstantiateFlagDeviceLaunch).
Uploaded to the device before any device-side launch, either explicitly via cudaGraphUpload(exec, stream) or implicitly on a first host launch. A device launch with no prior upload errors out.
Composed only of kernel, memcpy, memset, and child-graph nodes (the device-launchable node set).

Device graphs cannot be launched into ordinary streams; they use distinct named stream constants that select the mode. (NVIDIA CUDA Programming Guide, CUDA Graphs; NVIDIA Technical Blog, Device Graph Launch)

Host setup:

cudaGraphExec_t exec;
// Instantiate for device launch.
cudaGraphInstantiate(&exec, graph, cudaGraphInstantiateFlagDeviceLaunch);
// Upload device resources before the kernel that will relaunch it runs.
cudaGraphUpload(exec, stream);

scheduler<<<1, 1, 0, stream>>>(exec, /* state */ d_state);
cudaStreamSynchronize(stream);

Device-side launch from a kernel:

__global__ void scheduler(cudaGraphExec_t work, int* state) {
    // Fire-and-forget: dispatch immediately, independent of this grid.
    cudaGraphLaunch(work, cudaStreamGraphFireAndForget);

    if (should_continue(state)) {
        // Self-relaunch via tail launch builds a GPU-resident loop:
        // this graph runs again after the current execution completes.
        cudaGraphLaunch(cudaGetCurrentGraphExec(), cudaStreamGraphTailLaunch);
    }
}

Device-side named stream constants and helper:

cudaStreamGraphFireAndForget: graph runs immediately, independent of the launching graph and of other fire-and-forget launches.
cudaStreamGraphTailLaunch: graph runs after the launching graph (and its fire-and-forget launches) complete.
cudaStreamGraphFireAndForgetAsSibling: fire-and-forget enqueued as a sibling of the launching graph rather than as its child.
cudaGetCurrentGraphExec(): returns the currently executing graph's handle, used for the self-relaunch (tail-launch-to-self) loop.

(NVIDIA Technical Blog, Device Graph Launch; NVIDIA CUDA Runtime API, Graph Management)

Pending-launch limits (do not exceed)¶

Device launch enforces hard pending-launch limits per execution. The sibling CUDA graphs page records the documented caps: up to 120 total fire-and-forget graphs per execution, up to 255 pending tail launches, and only one pending self-relaunch at a time. Build scheduler loops around the single-self-relaunch rule rather than queueing many. (NVIDIA Technical Blog, Device Graph Launch; Fregly, Ch. 12)

This standard-library validator checks an application's pending-launch budget before graph construction. It does not execute CUDA or validate graph nodes.

LIMITS = {
    "fire_and_forget": 120,
    "pending_tail": 255,
    "self_relaunch": 1,
}


def validate_launch_budget(**requested):
    unknown = set(requested) - set(LIMITS)
    if unknown:
        raise ValueError(f"unknown launch classes: {sorted(unknown)}")
    for name, limit in LIMITS.items():
        value = requested.get(name, 0)
        if not isinstance(value, int) or value < 0:
            raise ValueError(f"{name} must be a non-negative integer")
        if value > limit:
            raise ValueError(f"{name}={value} exceeds documented limit {limit}")
    return True


assert validate_launch_budget(
    fire_and_forget=120, pending_tail=255, self_relaunch=1)

for invalid in (
    {"fire_and_forget": 121},
    {"pending_tail": 256},
    {"self_relaunch": 2},
    {"pending_tail": -1},
):
    try:
        validate_launch_budget(**invalid)
    except ValueError:
        pass
    else:
        raise AssertionError(f"invalid budget accepted: {invalid}")

print("device-launch budget validation: all asserts passed")

Executed output:

device-launch budget validation: all asserts passed

How to maintain it¶

Decide CDP vs device graph by structure: variable, irregular, possibly recursive child work fits CDP; a fixed-structure DAG replayed with runtime data fits device graph launch.
Watch device-launch overhead. Each child/graph launch consumes a pending-launch slot and pool memory; over-decomposition regresses throughput. Compare against a persistent kernel pulling from a work queue before committing to fine-grained device launches. (Fregly, Ch. 12)
Profile with a Nsight workflow: confirm the host-side per-decision gaps actually collapse and that device-launch overhead does not eat the savings. Do not assume the illustrative ~2x; measure on your target GPU.

How to run it in production¶

Bound every data-dependent launch count, expose rejected-work and pending-launch metrics, and keep a host-launched or persistent-kernel fallback. Validate uploaded device graphs during process warm-up. Compare outputs with the fallback for empty work, one item, maximum configured fan-out, and an over-limit request.

Reference templates only. APIs, flags, and limits are grounded in the cited book chapter and official NVIDIA docs. Not hardware-tested here. Benchmark on your target before relying on any figure.

Programmatic Dependent Launch (PDL): overlapping host-launched kernels¶

PDL solves a narrower, different problem than the two device-initiated mechanisms above: two kernels that are still both launched by the host, back to back in the same stream, but where the second kernel has some prologue work (loading constants, zeroing an output buffer, indexing setup) that does not actually depend on the first kernel's result. Ordinarily the second kernel's grid cannot start until every block of the first kernel has exited, even though only part of the second kernel's work needs the first kernel's output. PDL, available starting with compute capability 9.0, lets the second (secondary) kernel's thread blocks begin executing before the first (primary) kernel's grid has fully completed, so that non-dependent prologue overlaps the primary's tail instead of waiting behind it.¹

This is not CDP and not device graph launch: no kernel launches anything from device code here. Both kernels are still enqueued by the host in program order; PDL only relaxes when the second one is allowed to start, not who initiates it. Use the three mechanisms this way:

PDL: two host-launched kernels in one stream, one has non-dependent prologue work worth overlapping with the other's tail. No device-side launch involved.
CDP: a kernel must launch a new kernel whose existence or shape is only known from data the GPU just computed.
Device graph launch: a kernel must launch a previously built, fixed-structure graph, based on runtime data, without a host round trip.

The trigger/synchronize pair¶

The primary kernel opts in to being interruptible by calling cudaTriggerProgrammaticLaunchCompletion() from every thread block once it has produced everything the secondary kernel will need; the secondary kernel calls cudaGridDependencySynchronize() before it touches anything the primary produced.¹ If the primary never calls the trigger, it "implicitly occurs after all thread blocks in the primary kernel exit," so correctness never depends on the primary remembering to call it, only performance does.¹

// Primary kernel: signal as soon as everything the secondary needs is written.
__global__ void primary_kernel(float* out, int n) {
    // ... compute and write `out` ...
    cudaTriggerProgrammaticLaunchCompletion();   // secondary MAY start now
    // any remaining primary-only work continues here, potentially overlapping
    // with the secondary kernel's already-running prologue
}

// Secondary kernel: do non-dependent setup first, then synchronize before
// touching anything the primary produced.
__global__ void secondary_kernel(const float* in, float* result, int n) {
    // Non-dependent prologue: safe to run before the primary finishes.
    // e.g. zero a local accumulator, load constants, compute this thread's index.
    cudaGridDependencySynchronize();             // blocks until primary's global writes are visible
    // From here on, `in` (written by the primary) is safe to read.
}

The secondary kernel must be launched with the extensible launch API and an explicit opt-in attribute; a plain triple-chevron launch does not enable PDL:

cudaLaunchAttribute attribute[1];
attribute[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attribute[0].val.programmaticStreamSerializationAllowed = 1;

cudaLaunchConfig_t config_secondary = {};
config_secondary.gridDim = grid;
config_secondary.blockDim = block;
config_secondary.attrs = attribute;
config_secondary.numAttrs = 1;

cudaLaunchKernelEx(&config_secondary, secondary_kernel, in, result, n);

Memory visibility, and why `cudaGridDependencySynchronize()` is not optional¶

Without synchronizing, "the secondary thread blocks might launch before data written by the primary kernel is visible."¹ cudaGridDependencySynchronize() is the standard mechanism used here to establish that visibility; NVIDIA also permits other synchronization mechanisms that provide the required ordering. The trigger call alone only permits the secondary to start. A secondary kernel that reads primary output without any valid dependency synchronization has a data race that can appear to work in testing and corrupt results intermittently in production.

Opportunistic overlap, and the deadlock risk of relying on it¶

NVIDIA is explicit that the overlap PDL enables is "opportunistic and not guaranteed to lead to concurrent kernel execution," and separately warns: "reliance on concurrent execution in this manner is unsafe and can lead to deadlock."¹ Concretely: if the secondary kernel's non-dependent prologue itself waits on something only the primary kernel (or a third kernel) can supply, such as a lock, a semaphore, or a cross-grid barrier, and the runtime happens not to schedule the primary concurrently with that wait, the wait never resolves. Treat PDL purely as a latency-hiding optimization for independent prologue work, never as a substitute for correctness-critical ordering, and never write a secondary-kernel prologue that blocks on anything besides cudaGridDependencySynchronize() itself.

PDL inside CUDA Graphs¶

The same trigger/synchronize relationship can be expressed as a graph edge instead of a raw stream launch, via cudaGraphDependencyTypeProgrammatic on the edge between the primary and secondary kernel nodes, so a graph capturing this pattern replays with the same opportunistic overlap on replay.¹ This graph edge type is unrelated to device graph launch above: it changes how two nodes inside one graph overlap, not who launches the graph.

When to reach for PDL (and when not)¶

Use PDL when profiling shows a secondary kernel's launch is fully serialized behind a primary kernel it does not fully depend on, and the secondary has real non-dependent work at its start worth overlapping (not just a few instructions). Skip it when the secondary kernel's very first instruction already needs the primary's output: there is nothing to overlap, and the trigger/synchronize plumbing is pure overhead. Skip it, too, if you cannot audit the secondary kernel's prologue for hidden dependencies on the primary or on a third party, given the deadlock warning above.

Failure modes¶

A CDP2 parent reads child output before a tail-launch boundary.
A device graph is launched before upload or contains unsupported node types.
Pending fire-and-forget, tail, or self-relaunch counts exceed documented limits.
Tiny child grids cost more to launch than a persistent work queue.
PDL code reads primary output before valid dependency synchronization or assumes overlap is guaranteed.

References¶

Chris Fregly, AI Systems Performance Engineering (O'Reilly), Chapter 12: "Dynamic Scheduling, CUDA Graphs, and Device-Initiated Kernel Orchestration" (CUDA Dynamic Parallelism, device-initiated and self-relaunching graph launch, pending-launch limits, ~2x device-launch latency figure).
NVIDIA, CUDA C++ Programming Guide — Dynamic Parallelism: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/dynamic-parallelism.html
NVIDIA, CUDA C++ Programming Guide — CUDA Graphs (Device Graph Launch): https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cuda-graphs.html
NVIDIA Technical Blog, Enabling Dynamic Control Flow in CUDA Graphs with Device Graph Launch: https://developer.nvidia.com/blog/enabling-dynamic-control-flow-in-cuda-graphs-with-device-graph-launch/
NVIDIA, CUDA Runtime API — Graph Management: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__GRAPH.html
NVIDIA Developer Forums, cudaDeviceSynchronize from device code is deprecated: https://forums.developer.nvidia.com/t/cudadevicesynchronize-from-device-code-is-deprecated/215900
NVIDIA, CUDA C++ Programming Guide — Programmatic Dependent Launch and Synchronization: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/programmatic-dependent-launch.html

NVIDIA CUDA Programming Guide, "Programmatic Dependent Launch and Synchronization": available from compute capability 9.0; primary calls cudaTriggerProgrammaticLaunchCompletion() from every thread block, which "implicitly occurs after all thread blocks in the primary kernel exit" if never called explicitly; secondary calls cudaGridDependencySynchronize() before using primary output, since "the secondary thread blocks might launch before data written by the primary kernel is visible"; secondary launch requires the extensible launch API with cudaLaunchAttributeProgrammaticStreamSerialization set via cudaLaunchKernelEx; overlap is "opportunistic and not guaranteed to lead to concurrent kernel execution," and "reliance on concurrent execution in this manner is unsafe and can lead to deadlock"; the same relationship is expressible as a graph edge via cudaGraphDependencyTypeProgrammatic. https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/programmatic-dependent-launch.html ↩↩↩↩↩↩