Lanzamiento de kernels desde un kernel
Recursión y refinamiento en el dispositivo
Lanzamiento de kernels desde un kernel es una lección gratuita de CUDA Academy en CoddyKit. Esta es la lección 1 de 4. Puedes leer la lección completa abajo gratuitamente — luego la practicas en el navegador con un editor de código integrado y un tutor de IA 24/7. Forma parte de la ruta de aprendizaje de CUDA Academy, y tu progreso se sincroniza en la web y la app de CoddyKit. El curso de CUDA Academy incluye 4 lecciones en total.
Partes de esta lección aún no han sido traducidas y se muestran en inglés.
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! 🎉
Preguntas frecuentes
¿La lección «Lanzamiento de kernels desde un kernel» es gratis?
Sí — el texto completo de «Lanzamiento de kernels desde un kernel» es gratis para leer aquí en la web. Para practicarla de forma interactiva (editor de código integrado y tutor de IA 24/7) y desbloquear el resto del curso de CUDA Academy, actualiza a CoddyKit PRO. El curso de CUDA Academy incluye 4 lecciones en total.
¿Qué aprenderé en «Lanzamiento de kernels desde un kernel»?
Recursión y refinamiento en el dispositivo Practicas CUDA Academy con código real que ejecutas directamente en el navegador, y un tutor de IA 24/7 responde tus preguntas mientras trabajas en la lección.
¿Necesito experiencia previa para empezar CUDA Academy?
No se requiere experiencia previa. CUDA Academy en CoddyKit está estructurado para principiantes hasta estudiantes avanzados, así que puedes empezar aquí o desde el inicio y avanzar a tu ritmo. Esta es la lección 1 de 4.
¿Cuánto tiempo toma la lección «Lanzamiento de kernels desde un kernel»?
La mayoría de las lecciones de CoddyKit toman alrededor de 5–10 minutos. Cada una es compacta e interactiva, así que avanzas constantemente y retomas exactamente por donde dejaste en la web y la app.
¿Puedo escribir y ejecutar código en esta lección de CUDA Academy?
Sí. Cada lección de CUDA Academy incluye un editor de código integrado, así que escribes y ejecutas código real directamente en tu navegador y obtienes retroalimentación instantánea de IA — sin configuración local necesaria.
Todas las lecciones de este curso
- Lanzamiento de kernels desde un kernel
- Cuándo compensa el paralelismo dinámico
- Captura del trabajo en un grafo
- Repetición de grafos para reducir la sobrecarga