CUDA Dynamic Parallelism, in depth: CDP2 tail launch and fire-and-forget streams, memory visibility, launch-pool limits and migrating from CDP1

The execution model

A kernel launched from the host is a parent grid; any thread in it can launch a child grid with the usual <<<grid, block, smem, stream>>> syntax, and children can launch grandchildren. Nesting is strict: a parent grid is not complete until every grid launched by its threads has completed, and if all its threads exit early, an implicit synchronization waits for the children. The host sees one launch that finishes when the whole tree does.

Launch configuration is inherited. Children get the device's cache configuration and limits, such as stack size, and per-kernel attributes set from the host still apply when the kernel is launched from the device; the device cannot change them. There is no multi-GPU support from device code. The legacy CDP1 documentation puts the maximum nesting depth at 24, but memory for each level is the practical limit, so design for a depth of a few levels with an explicit cap in your own code.

Building requires relocatable device code and the device runtime library:

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

# CMake
set_target_properties(spmv PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
target_link_libraries(spmv PRIVATE CUDA::cudadevrt)

Linking the device runtime has a cost even for kernels that never launch anything: the programming guide warns that the runtime's tracking software can slow any kernel running while dynamic launches are being managed. Keep CDP in the binaries that need it.

Streams on the device

Device code has four kinds of stream, and choosing among them is most of the design work.

Device streams must be created non-blocking; cudaStreamCreate does not compile for the device, and neither cudaStreamSynchronize nor cudaStreamQuery exists there. Events work only for ordering between streams: create them with cudaEventDisableTiming and use cudaStreamWaitEvent; timing and host-style event synchronization are not available. Neither named stream can record or wait on events.

Tail launches from one grid run one after another in launch order. To fan out after a parent completes, tail-launch a single small kernel that launches the concurrent children with fire-and-forget. Tail launch also inserts itself before the next grid in the parent's host stream, so host-side ordering keeps working.

Memory visibility between parent and child

Parent and child share global, constant and mapped memory, with weak consistency. There is exactly one point at which the child's view is guaranteed to match the parent thread's: the moment of launch. Everything the launching thread wrote before the launch is visible to the child. If other threads in the block wrote data the child needs, put a __syncthreads() before the launch.

In the other direction, with cudaDeviceSynchronize() gone, the parent grid can never safely read what its children wrote. The only way to consume child results before returning to the host is a kernel in cudaStreamTailLaunch.

Shared and local memory are private. Passing a pointer to a parent's __shared__ array or to a local variable is undefined behaviour, and the compiler cannot always tell when a variable has spilled to local memory. Pass only pointers from cudaMalloc, device new or __device__ globals; __isGlobal(ptr) checks at runtime.

Worked example: SpMV with heavy-row children

Graph neural network aggregation and sparse attention both reduce to sparse matrix-vector products over power-law matrices: most rows have a handful of nonzeros, and a few hub rows have hundreds of thousands. One thread per row leaves the warp holding a hub row running long after its neighbours finish. With CDP the light rows stay one-thread-per-row, and each heavy row launches its own child grid. A tail launch then applies the activation once every row is complete.

#define HEAVY 4096           // nonzeros above which a row gets its own grid

__global__ void heavy_row(const int* colind, const float* val, const float* x,
                          float* y, int start, int end, int row) {
    float acc = 0.f;
    for (int i = start + blockIdx.x * blockDim.x + threadIdx.x; i < end;
         i += gridDim.x * blockDim.x)
        acc += val[i] * x[colind[i]];
    for (int o = 16; o > 0; o >>= 1) acc += __shfl_down_sync(0xffffffff, acc, o);
    if ((threadIdx.x & 31) == 0) atomicAdd(&y[row], acc);   // y[row] zeroed by parent
}

__global__ void relu_finish(float* y, int n) {               // runs after every child
    int r = blockIdx.x * blockDim.x + threadIdx.x;
    if (r < n) y[r] = fmaxf(y[r], 0.f);
}

__global__ void spmv_parent(const int* rowptr, const int* colind, const float* val,
                            const float* x, float* y, int n, int* launch_failures) {
    int r = blockIdx.x * blockDim.x + threadIdx.x;
    if (r < n) {
        int s = rowptr[r], e = rowptr[r + 1];
        if (e - s <= HEAVY) {
            float acc = 0.f;
            for (int i = s; i < e; ++i) acc += val[i] * x[colind[i]];
            y[r] = acc;
        } else {
            y[r] = 0.f;                                      // visible to the child
            int blocks = min((e - s + 255) / 256, 1024);
            heavy_row<<<blocks, 256, 0, cudaStreamFireAndForget>>>(colind, val, x, y, s, e, r);
            if (cudaGetLastError() != cudaSuccess) atomicAdd(launch_failures, 1);
        }
    }
    if (blockIdx.x == 0 && threadIdx.x == 0)
        relu_finish<<<(n + 255) / 256, 256, 0, cudaStreamTailLaunch>>>(y, n);
}

Walk through it. Each parent thread owns one row. A light row is computed inline. A heavy row zeroes its output and launches heavy_row with fire-and-forget, because heavy rows are independent and nothing in the parent needs their result. Thread 0 of block 0 issues exactly one tail launch; it does not run until every block of the parent and every fire-and-forget child has finished, so relu_finish sees final values. The host launches spmv_parent once and then reads launch_failures.

Two details matter. First, the launch result is checked with cudaGetLastError() in the launching thread; a failed device launch is otherwise silent, and that row would quietly read zero. Second, the number of heavy rows is bounded by the launch pool, which the next section covers. If a matrix can have more heavy rows than slots, raise the limit or batch heavy rows into one child that processes a list.

Test the kernel against a reference such as a CPU loop or a library SpMV on matrices with known hub rows. Expect small differences: the child grid accumulates with atomicAdd, so the order of floating-point additions in heavy rows changes from run to run, and results are not bitwise reproducible. If you need determinism, have each child block write a partial sum to a scratch array and let the tail-launched kernel add the partials in a fixed order. Tune HEAVY by measurement: too low and launch overhead dominates, too high and the long rows return to stalling their warps.

The launch pool and other limits

Every pending device launch occupies a slot in a fixed-size pool. The default is 2048 slots, and event slots are twice the launch limit. When the pool is full, a device launch fails with cudaErrorLaunchOutOfResources and an event allocation with cudaErrorMemoryAllocation. CDP1 had a slower virtualized overflow pool; CDP2 does not, so you must size the limit yourself, from the host, before the first kernel runs:

size_t heavy = count_heavy_rows(host_rowptr, n, HEAVY);      // known before launch
size_t slots = heavy + 64;                                   // + tail and headroom
cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, std::max(slots, (size_t)2048));

If you cannot bound the count, change the design: collect heavy rows into a global work list with an atomic counter, and tail-launch one kernel that processes the whole list. That turns n launches into one and removes the pool from the failure surface.

Migrating CDP1 code

CDP1 code typically launched a child, called cudaDeviceSynchronize(), and continued with the results. Under CDP2 that does not compile, and code built with an older toolkit that references it fails to load on compute capability 9.0 devices with cudaErrorSymbolNotFound. The mechanical fix is to split the parent at each synchronization point:

1. Move everything after the sync into a continuation kernel.
2. Move the state the parent kept in registers or shared memory into global memory, since the continuation is a new grid.
3. Tail-launch the continuation from one thread, after the children have been launched.
4. If the old code synchronized inside a loop, move the loop out of the kernel: drive it from the host, or build it as a CUDA graph with a while conditional node, or as a device graph that tail-launches itself, which the guide documents for device graphs (one self-launch pending at a time).

Defining CUDA_FORCE_CDP1_IF_SUPPORTED keeps CDP1 working below compute capability 9.0 only; it is a stopgap, not a migration.

Failure modes

• Spin-waiting on a child. The guide makes no guarantee that a child starts before its parent reaches an implicit synchronization point. A parent polling a flag the child sets can deadlock. Use tail launch.
• Silent launch failures. Pool exhaustion or a bad configuration fails a single device launch, and nothing reaches the host unless you record it.
• Shared or local pointers passed to children, including through cudaMemcpyAsync on the device, which may itself launch a kernel.
• Reading child output in the parent: it may appear to work in testing and then fail under load.
• Launch storms. One child per thread for tiny amounts of work costs more in launch overhead than it saves. Launch only above a threshold, as with HEAVY above.
• Stream handles crossing grids. A stream created in one grid used in a child is undefined.

Trade-offs and alternatives

A rule of thumb: if the irregularity has one or two levels and the host can afford one synchronization, the two-pass design is easier to maintain. Choose CDP when recursion depth or fan-out is data-dependent across several levels. For overlapping host-launched work see multi-stream execution and stream capture; for whether the parent grid leaves room for children see occupancy and warp scheduling.

What to do next

1. Grep your code for device-side cudaDeviceSynchronize and plan tail-launch rewrites before moving to compute capability 9.0 hardware.
2. Build with -rdc=true and -lcudadevrt only in the targets that launch from the device.
3. Check cudaGetLastError() after every device launch and count failures in global memory.
4. Compute an upper bound on pending launches and set cudaLimitDevRuntimePendingLaunchCount before the first kernel.
5. Use fire-and-forget for independent children and exactly one tail launch for the continuation.
6. Profile against a host two-pass version of the same algorithm and keep whichever is faster and simpler.