커널에서 커널 실행하기
장치 측 재귀와 세분화를 살펴봅니다
커널에서 커널 실행하기은(는) CoddyKit의 무료 CUDA Academy 강의입니다. 이것은 4개 중 1번째 강의입니다. 아래에서 전체 강의를 무료로 읽을 수 있으며, 내장 코드 에디터와 24/7 AI 튜터와 함께 브라우저에서 직접 실습할 수 있습니다. 이 강의는 CUDA Academy 학습 경로의 일부이며, 진행 상황이 웹과 CoddyKit 앱에 동기화됩니다. CUDA Academy 강의에는 총 4개의 강의가 포함되어 있습니다.
이 강의의 일부는 아직 번역되지 않았으며 영어로 표시됩니다.
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! 🎉
자주 묻는 질문
“커널에서 커널 실행하기” 강의는 무료인가요?
네 — “커널에서 커널 실행하기” 전체 내용을 이 웹사이트에서 무료로 읽을 수 있습니다. 인터랙티브하게 실습하려면(내장 코드 에디터와 24/7 AI 튜터), CoddyKit PRO로 업그레이드하면 CUDA Academy 강의 전체를 잠금 해제할 수 있습니다. CUDA Academy 강의에는 총 4개의 강의가 포함되어 있습니다.
“커널에서 커널 실행하기”에서 뭘 배우나요?
장치 측 재귀와 세분화를 살펴봅니다 브라우저에서 직접 실행하는 실습 코드로 CUDA Academy을(를) 배우며, 24/7 AI 튜터가 강의를 진행하면서 질문에 답변해줍니다.
CUDA Academy을(를) 시작하는 데 경험이 필요한가요?
사전 경험은 필요하지 않습니다. CoddyKit의 CUDA Academy은(는) 초급자부터 고급 학습자까지를 위해 구성되어 있으므로, 여기서 시작하거나 처음부터 시작할 수 있으며 자신의 속도대로 진행할 수 있습니다. 이것은 4개 중 1번째 강의입니다.
“커널에서 커널 실행하기” 강의는 얼마나 걸리나요?
대부분의 CoddyKit 강의는 약 5~10분이 소요됩니다. 각 강의는 간결하고 인터랙티브하여 꾸준한 진행이 가능하며, 웹과 앱에서 중단한 부분부터 바로 시작할 수 있습니다.
이 CUDA Academy 강의에서 코드를 작성하고 실행할 수 있나요?
네. 모든 CUDA Academy 강의에는 내장 코드 에디터가 포함되어 있으므로, 브라우저에서 바로 실제 코드를 작성하고 실행한 후 즉시 AI 피드백을 받을 수 있습니다 — 로컬 설정이 필요 없습니다.