Kernel aus einem Kernel starten
Rekursion und Verfeinerung auf dem Gerät
Kernel aus einem Kernel starten ist eine kostenlose CUDA Academy-Lektion auf CoddyKit. Dies ist Lektion 1 von 4. Du kannst die komplette Lektion unten kostenlos lesen – dann übst du sie direkt im Browser mit einem integrierten Code-Editor und einem KI-Tutor rund um die Uhr. Sie ist Teil des CUDA Academy-Lernpfads, und dein Fortschritt wird über Web und CoddyKit-App synchronisiert. Der CUDA Academy-Kurs umfasst insgesamt 4 Lektionen.
Teile dieser Lektion wurden noch nicht übersetzt und werden auf Englisch angezeigt.
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! 🎉
Häufig gestellte Fragen
Ist die Lektion „Kernel aus einem Kernel starten“ kostenlos?
Ja — der vollständige Text von „Kernel aus einem Kernel starten“ ist hier im Web kostenlos zu lesen. Um sie interaktiv zu üben (integrierter Code-Editor und 24/7 KI-Tutor) und den Rest des CUDA Academy-Kurses freizuschalten, upgrade auf CoddyKit PRO. Der CUDA Academy-Kurs umfasst insgesamt 4 Lektionen.
Was lerne ich in „Kernel aus einem Kernel starten“?
Rekursion und Verfeinerung auf dem Gerät Du übst CUDA Academy mit praktischem Code, den du direkt im Browser ausführst, und ein 24/7 KI-Tutor beantwortet deine Fragen während du die Lektion bearbeitest.
Brauche ich Erfahrung, um CUDA Academy zu starten?
Keine Vorkenntnisse erforderlich. CUDA Academy auf CoddyKit ist für Anfänger bis fortgeschrittene Lernende strukturiert, sodass du hier starten oder von Anfang an beginnen und in deinem eigenen Tempo voranschreiten kannst. Dies ist Lektion 1 von 4.
Wie lange dauert die Lektion „Kernel aus einem Kernel starten“?
Die meisten CoddyKit-Lektionen dauern etwa 5–10 Minuten. Jede ist kompakt und interaktiv, sodass du stetig Fortschritte machst und genau dort weitermachst, wo du aufgehört hast – im Web und in der App.
Kann ich in dieser CUDA Academy-Lektion Code schreiben und ausführen?
Ja. Jede CUDA Academy-Lektion enthält einen integrierten Code-Editor, sodass du echten Code direkt in deinem Browser schreibst und ausführst und sofort KI-Feedback erhältst — ohne lokale Einrichtung erforderlich.
Alle Lektionen in diesem Kurs
- Kernel aus einem Kernel starten
- Wann sich Dynamic Parallelism lohnt
- Arbeit in einem Graphen erfassen
- Graphen zur Reduzierung des Overheads wiedergeben