Lançando kernels a partir de um kernel
Recursão e refinamento no dispositivo.
Lançando kernels a partir de um kernel é uma aula grátis de CUDA Academy no CoddyKit. Esta é a aula 1 de 4. Você pode ler a aula completa abaixo gratuitamente — depois pratica ao vivo no navegador com um editor de código integrado e um tutor de IA 24/7. Faz parte do caminho de aprendizado de CUDA Academy, e seu progresso é sincronizado entre a web e o app CoddyKit. O curso de CUDA Academy inclui 4 aulas no total.
Partes desta aula ainda não foram traduzidas e aparecem em 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! 🎉
Perguntas Frequentes
A aula “Lançando kernels a partir de um kernel” é grátis?
Sim — o texto completo de “Lançando kernels a partir de um kernel” é grátis para ler aqui na web. Para praticá-la interativamente (um editor de código integrado e um tutor de IA 24/7) e desbloquear o restante do curso de CUDA Academy, atualize para CoddyKit PRO. O curso de CUDA Academy inclui 4 aulas no total.
O que vou aprender em “Lançando kernels a partir de um kernel”?
Recursão e refinamento no dispositivo. Você pratica CUDA Academy com código prático que executa diretamente no navegador, e um tutor de IA 24/7 responde suas dúvidas enquanto trabalha na aula.
Preciso ter experiência prévia para começar CUDA Academy?
Nenhuma experiência prévia é necessária. CUDA Academy no CoddyKit é estruturado para alunos iniciantes até avançados, então você pode começar aqui ou desde o início e aprender no seu ritmo. Esta é a aula 1 de 4.
Quanto tempo leva a aula “Lançando kernels a partir de um kernel”?
A maioria das aulas CoddyKit leva cerca de 5–10 minutos. Cada uma é compacta e interativa, então você faz progresso constante e retoma exatamente de onde parou entre web e app.
Posso escrever e executar código nesta aula de CUDA Academy?
Sim. Cada aula de CUDA Academy inclui um editor de código integrado, então você escreve e executa código real direto no navegador e recebe feedback de IA instantaneamente — nenhuma configuração local necessária.
Todas as aulas deste curso
- Lançando kernels a partir de um kernel
- Quando o paralelismo dinâmico compensa
- Capturando trabalho em um grafo
- Reproduzindo grafos para reduzir sobrecarga