Launching Kernels from a Kernel
Device-side recursion and refinement.
Launching Kernels from a Kernel is a free CUDA Academy lesson on CoddyKit — lesson 1 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
Kernels Launching Kernels
With dynamic parallelism, a running GPU kernel can launch another kernel itself, no trip back to the CPU required. 🚀
Same Syntax, Device Side
The launch looks identical to a host launch: the familiar triple angle brackets work right inside device code.
__global__ void child(int* d);
__global__ void parent(int* d) {
child<<<1, 32>>>(d);
}Compute the Future Compute
A parent thread decides at runtime how much work is needed, then spawns a child kernel sized exactly for it.
child<<<blocks, threads>>>(buf);Built for Irregular Shapes
This shines when the work is data-dependent: a thread that finds more to do can launch more threads on the spot.
Device-Side Recursion
A kernel may even launch itself, giving you true recursion on the GPU for divide-and-conquer problems.
__global__ void solve(int lo, int hi) {
if (hi - lo > 1) solve<<<1, 2>>>(lo, mid);
}Parent and Child Streams
Each child launch joins a stream. By default children use the parent thread block stream, but you can name your own stream.
child<<<g, b, 0, myStream>>>(d);Synchronizing Inside a Kernel
A parent can wait for its children with cudaDeviceSynchronize called from device code, then read their results.
child<<<1, 64>>>(d);
cudaDeviceSynchronize();Implicit Child Completion
Even without a manual wait, all launched children are guaranteed complete before the parent kernel itself returns.
Memory Both Can See
Parent and child share global memory, so pointers passed down stay valid. Local and shared data, though, do not cross the launch.
Compile for Device Launch
You must compile with relocatable device code and link the device runtime so nested launches actually work.
nvcc -rdc=true -lcudadevrt prog.cuA Hard Launch Limit
Nesting is capped: there is a maximum launch depth, and going too deep returns an error instead of more kernels.
Quick Check
Where does a dynamically launched kernel originate?
Recap: Device Launches
A kernel can launch other kernels with the same triple-bracket syntax, even itself. Share global memory and compile with -rdc=true. Nice work! 🎉
Frequently asked questions
Is the “Launching Kernels from a Kernel” lesson free?
Yes — the full text of “Launching Kernels from a Kernel” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “Launching Kernels from a Kernel”?
Device-side recursion and refinement. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 1 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Launching Kernels from a Kernel” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- Launching Kernels from a Kernel
- When Dynamic Parallelism Pays
- Capturing Work into a Graph
- Replaying Graphs to Cut Overhead