Dynamic Parallelism

Kernels can launch other kernels from the device — useful for recursive/adaptive algorithms (quadtrees, adaptive mesh refinement, recursive sorts).

CUDA 12 introduced a revised model (CDP2): device-side cudaDeviceSynchronize() was removed; use tail launches (cudaStreamTailLaunch) and fire-and-forget streams. Launch overhead is significant — prefer flattening work or cooperative/persistent approaches when possible.

__global__ void child(int* data, int depth) { /* ... */ }

__global__ void parent(int* data, int depth) {
    if (threadIdx.x == 0 && depth < 4) {
        child<<<4, 64>>>(data, depth + 1);    // device-side launch
    }
}
// nvcc -rdc=true -arch=sm_80 dp.cu -lcudadevrt