cuda dynamic parallelism
**CUDA Dynamic Parallelism** is the **ability for GPU kernels to launch other GPU kernels directly from the device** — eliminating round-trips to the CPU for recursive or adaptive algorithms where the next work unit depends on computed results.
**Traditional GPU Programming Constraint**
- Old model: CPU → launch kernel → GPU runs → CPU reads results → CPU decides next work → launch next kernel.
- Round-trip CPU-GPU overhead: 10–50 μs per kernel launch.
- Problem: Algorithms needing recursive subdivision required hundreds of CPU-GPU round-trips.
**Dynamic Parallelism Solution**
```cuda
__global__ void parent_kernel(int* data, int n) {
if (n > THRESHOLD) {
// Launch child kernel from within GPU kernel
child_kernel<<>>(data, n/2);
cudaDeviceSynchronize(); // Wait for child
merge_results<<<1, 32>>>(data, n);
} else {
base_case(data, n);
}
}
```
- Child kernels: Inherit parent's CUDA context.
- Synchronization: `cudaDeviceSynchronize()` within kernel waits for all launched children.
- Stream inheritance: Children run on parent's stream by default.
**When Dynamic Parallelism Helps**
- **Adaptive mesh refinement**: Refine only high-error regions → launch child kernels for refined areas.
- **Quicksort on GPU**: Partition → recursively sort two halves from device.
- **Sparse BFS**: Expand only non-empty frontier — don't launch fixed-size kernels.
- **Traversal algorithms**: Octree, BVH traversal with unknown depth.
**Performance Considerations**
- Child launch overhead: ~500ns on modern NVIDIA GPUs (vs. 10-50μs CPU-to-GPU).
- Memory: Child grid descriptors stored in global memory — small overhead.
- Nesting: Up to 24 levels of nesting supported (CUDA 5.0+).
- Overhead vs. benefit: Only worthwhile when CPU launch overhead was the bottleneck.
**Alternatives**
- Persistent threads: One kernel with internal work queue instead of nested launches.
- CUDA Graphs: Pre-record dynamic work patterns if structure is known.
CUDA Dynamic Parallelism is **the key enabler for GPU-native recursive and adaptive algorithms** — it eliminates the synchronization bottleneck that forced CPU coordination for work-adaptive GPU programs, enabling fully GPU-resident implementations of tree algorithms and adaptive solvers.