The CUDA Handbook ▸ Chapter 7. Kernel Execution

7.4 Dynamic Parallelism

Dynamic parallelism, available since SM 3.5-class hardware, enables CUDA kernels to launch other CUDA kernels, and also to invoke various functions in the CUDA runtime. When using dynamic parallelism, a subset of the CUDA runtime (known as the device runtime) becomes available for use by threads running on the device.

CUDA 12 revised dynamic parallelism significantly. The original model—retroactively named CDP1—let a parent grid block on its children by calling cudaDeviceSynchronize() from device code. That device-side synchronization was expensive to implement and has been removed: CUDA 12 defaults to a second-generation model, CDP2, in which a kernel can launch children but cannot synchronize on them from device code. On GPUs of compute capability 9.0 (Hopper) and later, CDP2 is the only model available; on earlier hardware, the legacy CDP1 behavior can still be selected at compile time (Section 7.4.5). This section describes CDP2, noting differences from CDP1 where they matter.

Dynamic parallelism introduces the idea of “parent” and “child” grids. Any kernel invoked by another CUDA kernel (as opposed to host code, as done in all previous CUDA versions) is a “child kernel,” and the invoking grid is its “parent.” A parent grid is not considered complete until all of the child grids launched by its threads have completed; the runtime enforces this dependency implicitly.

Dynamic parallelism was designed to address applications that previously had to deliver results to the CPU, so the CPU could specify which work to perform on the GPU. Such “handshaking” disrupts CPU/GPU concurrency in the execution pipeline described in Section 2.5.1, in which the CPU produces commands for consumption by the GPU. The GPU’s time is too valuable for it to wait for the CPU to read and analyze results before issuing more work. Dynamic parallelism avoids these pipeline bubbles by enabling the GPU to launch work for itself from kernels.

Dynamic parallelism can improve performance in several cases:

• It enables initialization of data structures needed by a kernel before the kernel can begin execution. Previously, such initialization had to be taken care of in host code or by previously invoking a separate kernel.

• It enables simplified recursion for applications such as Barnes-Hut gravitational integration or hierarchical grid evaluation for aerodynamic simulations.

7.4.1 Scoping and Synchronization

With the notable exception of block and grid size, child grids inherit most kernel configuration parameters, such as the shared memory configuration (set by cudaDeviceSetCacheConfig()), from their parents.

Streams and events created by the device runtime are scoped to the grid that created them: they can only be used by that grid (they are not inherited by child grids), and they are automatically destroyed when the grid exits. (Under the legacy CDP1 model, the unit of scope was the thread block rather than the grid.)

CUDA guarantees that a parent grid is not considered complete until all of its children have finished. The parent and its children may execute concurrently, and the runtime may begin a child grid at any time after it is launched. Under CDP2, device code cannot force a child to run, nor block on its completion—the parent-side cudaDeviceSynchronize() that CDP1 provided for this purpose is no longer available (referencing it from device code is a compile error).

Because device code can no longer wait on a child grid, work that must run only after a child completes is launched rather than waited for. A kernel launched into the special cudaStreamTailLaunch stream is guaranteed to run only after the launching grid—and every grid that grid has itself launched—has completed. This “tail launch” is the CDP2 replacement for the parent-side cudaDeviceSynchronize() of CDP1: the continuation that would have followed the synchronize call is written as a separate kernel and scheduled into the tail-launch stream. A kernel launched into cudaStreamFireAndForget, by contrast, begins as soon as resources allow, with no ordering relative to the launching grid’s other children.

As on the host, operations within a given device stream are performed in the order of submission; operations can only execute concurrently if they are specified in different streams; and there is no guarantee that operations will, in fact, execute concurrently. If needed, synchronization primitives such as __syncthreads() can be used to coordinate the order in which a block’s threads submit work to a given stream.

Streams and events created on the device may not be used outside the grid that created them. Because there is no device-side operation that waits on pending child work, ordering across grids is expressed with tail launches, and any coordination among a block’s threads (for example, ensuring all have submitted their launches before a tail launch is issued) still relies on __syncthreads() or other block-level synchronization primitives (Section 8.6.2).

7.4.2 Memory Model

Parent and child grids share the same global and constant memory storage, but have distinct local and shared memory.

Global Memory

In CDP2, there is a single point at which a child grid’s view of global memory is fully consistent with the parent: the moment the child is invoked. All global memory operations performed by the parent thread before it launches the child are visible to the child grid.

The reverse direction is more restricted than it was under CDP1. Because the parent cannot synchronize on the child from device code, a child grid’s writes are not guaranteed to be visible to the parent grid during the parent’s own execution. Work that must observe the child’s results is instead placed in a kernel launched into the cudaStreamTailLaunch stream, which runs after the child—and all of the parent’s other launches—has completed.

Zero-copy memory has the same coherence and consistency guarantees as global memory.

Constant Memory

Constants are immutable and may not be modified from the device during kernel execution.

Taking the address of a constant memory object from within a kernel thread has the same semantics as for all CUDA programs¹², and passing that pointer between parents and their children is fully supported.

Shared and Local Memory

Shared and local memory is private to a thread block or thread, respectively, and is not visible or coherent between parent and child. When an object in one of these locations is referenced outside its scope, the behavior is undefined and would likely cause an error.

If nvcc detects an attempt to misuse a pointer to shared or local memory, it will issue a warning. Developers can use the __isGlobal() intrinsic to determine whether a given pointer references global memory.

Pointers to shared or local memory are not valid parameters to cudaMemcpy*Async() or cudaMemset*Async().

Local Memory

Local memory is private storage for an executing thread, and is not visible outside of that thread. It is illegal to pass a pointer to local memory as a launch argument when launching a child kernel. The result of dereferencing such a local memory pointer from a child will be undefined.

To guarantee that this rule is not inadvertently violated by the compiler, all storage passed to a child kernel should be allocated explicitly from the global memory heap.

Texture Memory

Concurrent accesses by parent and child may result in inconsistent data and should be avoided. That said, a degree of coherency between parent and child is enforced by the runtime.

A child kernel can use texturing to access memory written by its parent, but writes to memory by a child are not reflected in texture accesses by the parent grid; a kernel tail-launched after the child (Section 7.4.1) is needed to observe them.

Texture objects are well supported in the device runtime – they cannot be created or destroyed, but they can be passed in and used by any grid in the hierarchy (parent or child).

7.4.3 Streams and Events

Streams and events created by the device runtime can be used only within the grid that created them.

Alongside user-created streams, CDP2 provides two built-in named streams—cudaStreamFireAndForget and cudaStreamTailLaunch (Section 7.4.1)—plus the NULL stream. The NULL stream has different semantics in the device runtime than in the host runtime. On the host, synchronizing with the NULL stream forces a ‘join’ of all other streamed operations on the GPU (as described in Section 6.2.3); on the device, the NULL stream is its own stream, shared by the threads of a block, and any inter-stream synchronization must be performed using events.

When using the device runtime, streams must be created with the cudaStreamNonBlocking flag (a parameter to cudaStreamCreateWithFlags()).

The cudaStreamSynchronize() call is not supported; synchronization must be implemented in terms of events and cudaStreamWaitEvent().

Only the inter-stream synchronization capabilities of CUDA events are supported. As a consequence, cudaEventSynchronize(), cudaEventElapsedTime(), and cudaEventQuery() are not supported. Additionally, because timing is not supported, events must be created by passing the cudaEventDisableTiming flag to cudaEventCreateWithFlags().

7.4.4 Error Handling

Any function in the device runtime may return an error (cudaError_t). The error is recorded in a per-thread slot that can be queried by calling cudaGetLastError().

As with the host-based runtime, CUDA makes a distinction between errors that can be returned immediately (e.g., if an invalid parameter is passed to a memcpy function) and errors that must be reported asynchronously (e.g., if a launch performed an invalid memory access). If a child grid causes an error at runtime, CUDA will return an error to the host, not to the parent grid.

7.4.5 Compiling and Linking

Unlike the host runtime, developers must explicitly link against the device runtime’s static library when using the device runtime. On Windows, the device runtime is cudadevrt.lib; on Linux and MacOS, it is cudadevrt.a. When building with nvcc, this may be accomplished by appending -lcudadevrt to the command line, together with -rdc=true to enable the relocatable device code that dynamic parallelism requires.

By default, nvcc compiles for CDP2. On hardware of compute capability below 9.0, the legacy CDP1 model can be selected instead by compiling with -DCUDA_FORCE_CDP1_IF_SUPPORTED; CDP1 and CDP2 code cannot be mixed within a program, and compute capability 9.0 and later support CDP2 only.

7.4.6 Resource Management

Whenever a kernel launches a child grid, the child is considered a new nesting level, and the total number of levels is the nesting depth of the program. Launches may be nested to a hardware-defined maximum depth, and a launch that would exceed it fails. Because CDP2 has no device-side synchronization, the CDP1 notion of a synchronization depth—and the cudaLimitDevRuntimeSyncDepth limit that capped the depth at which cudaDeviceSynchronize() could be called—no longer applies.

Calling a device runtime function such as cudaMemcpyAsync() may invoke a kernel, increasing the nesting depth by 1.

The limits must be configured before the top-level kernel is launched from the host.

Because CDP2 never suspends a parent to wait on its children, it does not reserve backing store to save parent-grid state—eliminating the per-nesting-level reservation (up to 150 MB per level) that dominated the CDP1 footprint.

Memory Footprint

The device runtime system software reserves device memory for the following purposes:

• to track pending grid launches, and

• to serve as an allocation heap for malloc() and cudaMalloc() calls from kernels.

This memory is not available for use by the application, so some applications may wish to reduce the default allocations; and some applications may have to increase the default values in order to operate correctly. To change the default values, developers call cudaDeviceSetLimit(), as summarized in Table 7-3.

Pending Kernel Launches

When a kernel is launched, all associated configuration and parameter data is tracked until the kernel completes. This data is stored within a system-managed launch pool. The size of the launch pool is configurable by calling cudaDeviceSetLimit() from the host and specifying cudaLimitDevRuntimePendingLaunchCount.

Configuration Options

Resource allocation for the device runtime system software is controlled via the cudaDeviceSetLimit() API from the host program. Limits must be set before any kernel is launched, and may not be changed while the GPU is actively running programs.

Memory allocated by the device runtime must be freed by the device runtime. Also, memory is allocated by the device runtime out of a preallocated heap whose size is specified by the device limit cudaLimitMallocHeapSize.

The following named limits may be set:

Table 7-3. cudaDeviceSetLimit() values.

The legacy CDP1 model accepted a third limit, cudaLimitDevRuntimeSyncDepth, which set the maximum depth at which device-side cudaDeviceSynchronize() could be called (default 2). It has no effect under CDP2, which has no device-side synchronization.

7.4.7 Summary

Table 7-4 summarizes the key differences and limitations between the device runtime and the host runtime. Table 7-5 lists the subset of functions that may be called from the device runtime, along with any pertinent limitations.

Table 7-4. Device Runtime Limitations

Table 7-5. CUDA Device Runtime Functions