CUDA Academy · 课时

从内核启动内核

设备端递归与细化

从内核启动内核 是 CoddyKit 上的免费 CUDA Academy 课时。这是第 1 节课，共 4 节。你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 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.cu

A 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! 🎉

常见问题解答

「从内核启动内核」课时是免费的吗？

是的 — 「从内核启动内核」的完整文本可在网页上免费阅读。要进行交互式练习（内置代码编辑器和全天候 AI 导师）并解锁 CUDA Academy 课程的其余内容，请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。

「从内核启动内核」这节课中我会学到什么？

设备端递归与细化你通过在浏览器中直接运行的动手代码来练习 CUDA Academy，全天候 AI 导师会在你学习这节课的过程中回答你的问题。

学习 CUDA Academy 需要有经验吗？

无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者，你可以从这里开始或从头开始，按照自己的节奏学习。这是第 1 节课，共 4 节。

「从内核启动内核」课时需要多长时间？

大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动，所以你能稳步进步，并在网页和应用中从离开的地方继续。

我能在这节 CUDA Academy 课中编写并运行代码吗？

能。每节 CUDA Academy 课都包含内置代码编辑器，你可以在浏览器中直接编写并运行真实代码，并获得即时 AI 反馈 — 无需本地设置。

此课程中的所有课时

← 返回 CUDA Academy