3、第一个动态并行程序：在GPU上启动内核（cudaLaunchDevice），父子内核概念
好，咱们直接动手。这一章的目标很明确——写一个能在GPU上自己启动自己的程序。说白了，就是让内核函数里再调一个内核函数。这在CUDA里叫动态并行（Dynamic Parallelism）。
我记得第一次接触这个概念时，心里还挺激动的。以前写CUDA程序，所有内核都得从CPU端发射，像放烟花一样，你点一次火，它炸一次。现在好了，GPU自己就能点火，想炸几次炸几次。嗯，这感觉就像从遥控车升级到了自动驾驶。
3.1 什么是父子内核？
先理清概念。动态并行里有两个角色：
- 父内核（Parent Kernel）：从CPU端启动的那个内核。它是整个并行任务的起点。
- 子内核（Child Kernel）：在父内核内部，通过
cudaLaunchDevice启动的新内核。子内核还可以再启动子内核，形成嵌套。
我习惯把父内核想象成「包工头」，子内核就是「工人」。包工头接到活（CPU下发任务），然后分派给工人去干。工人干不完还可以再招临时工（子内核的子内核）。
关键点：子内核的执行是异步的。父内核启动子内核后，不会等它执行完。父内核继续往下走，子内核在另一个执行流里跑。这一点和CPU端的cudaLaunchKernel行为一致。
3.2 第一个动态并行程序：从父内核启动子内核
咱们写一个最简单的例子。父内核里启动一个子内核，子内核只打印一句话。注意，这里有个坑——GPU上的printf在动态并行里行为有点特殊，我后面会讲。
先看代码：
// 子内核：简单打印
__global__ void childKernel(int *data) {
printf("我是子内核，线程ID: %d\n", threadIdx.x);
// 做一些计算
data[threadIdx.x] = threadIdx.x * 2;
}
// 父内核：启动子内核
__global__ void parentKernel(int *data) {
printf("我是父内核，线程ID: %d，准备启动子内核\n", threadIdx.x);
// 配置子内核的启动参数
dim3 gridDim(1); // 子内核的网格大小
dim3 blockDim(256); // 子内核的块大小
void *args[] = { &data }; // 参数列表
// 启动子内核
cudaLaunchDevice((void*)childKernel, gridDim, blockDim, args, 0, NULL);
printf("父内核：子内核已启动，继续执行\n");
}
int main() {
int *d_data;
cudaMalloc(&d_data, 256 * sizeof(int));
// 启动父内核
parentKernel<<<1, 1>>>(d_data);
cudaDeviceSynchronize();
printf("所有内核执行完毕\n");
cudaFree(d_data);
return 0;
}这段代码里，父内核只有一个线程（<<<1, 1>>>），它启动了一个有256个线程的子内核。你想想看，如果不用动态并行，你得在CPU端写两次cudaLaunchKernel。现在一次搞定。
我的习惯：写动态并行代码时，我总会在父内核和子内核里都加上printf，方便调试。但要注意，GPU上的printf输出是缓冲的，而且动态并行里的输出顺序可能和你想的不一样。我曾经因为这个排查了半天，后来发现是缓冲区没刷新。
3.3 cudaLaunchDevice 的细节
cudaLaunchDevice是动态并行的核心API。它的签名长这样：
cudaError_t cudaLaunchDevice(void *func, dim3 gridDim, dim3 blockDim,
void **args, size_t sharedMem, cudaStream_t stream);参数解释：
| 参数 | 说明 |
|---|---|
func |
子内核的函数指针，需要强转成void* |
gridDim |
子内核的网格维度 |
blockDim |
子内核的块维度 |
args |
参数指针数组，每个元素指向一个参数 |
sharedMem |
子内核的动态共享内存大小，单位字节 |
stream |
子内核运行的流，传NULL表示默认流 |
这里有个容易踩的坑：args参数。它必须是一个指针数组，每个元素指向实际参数。比如上面的例子，data本身已经是一个设备指针，所以&data就是指向这个指针的指针。嗯，有点绕，但写多了就习惯了。
我曾经犯过的错：直接把data传进去，忘了取地址。结果子内核拿到的参数全是垃圾值。排查了整整一个下午。记住：args里放的是参数的地址，不是参数本身。
3.4 动态并行的限制与注意事项
动态并行不是万能的。我总结了几条铁律：
- 计算能力要求：只有SM 3.5及以上（Kepler之后）的GPU才支持。你可以在代码里用
cudaGetDeviceProperties检查。 - 嵌套深度：最大嵌套深度是24层（包括父内核）。别问我怎么知道的，我曾经试图递归到100层，然后GPU直接挂了。
- 同步问题：父内核和子内核之间没有隐式同步。如果你需要子内核的结果，得用
cudaDeviceSynchronize或者__syncthreads。 - 内存限制：每个线程启动子内核时，会消耗一定的本地内存和寄存器。启动太多子内核可能导致栈溢出。
说白了，动态并行适合那些「任务量不确定」的场景。比如八叉树构建、自适应网格细化、稀疏矩阵操作。如果任务量是固定的，老老实实用CPU端启动更靠谱。
3.5 实战：用动态并行实现递归求和
咱们来点实际的。写一个递归归约求和，用动态并行实现。父内核把数据分成两半，启动两个子内核分别求和，然后合并结果。
__global__ void reduceKernel(float *data, int n, float *result) {
if (n == 1) {
*result = data[0];
return;
}
int half = n / 2;
float *left_result, *right_result;
// 分配临时空间（实际项目中应该预分配）
// 这里简化处理
// 启动左半部分子内核
dim3 grid(1);
dim3 block(1);
void *args_left[] = { &data, &half, &left_result };
cudaLaunchDevice((void*)reduceKernel, grid, block, args_left, 0, NULL);
// 启动右半部分子内核
int right_start = half;
void *args_right[] = { &data + right_start, &half, &right_result };
cudaLaunchDevice((void*)reduceKernel, grid, block, args_right, 0, NULL);
// 等待子内核完成（需要同步机制）
// 这里省略同步代码，实际需要cudaDeviceSynchronize
*result = *left_result + *right_result;
}这个例子展示了动态并行的核心思想：把大问题分解成小问题，每个小问题再递归分解。不过要注意，上面的代码为了简洁省略了同步和内存管理。实际项目中，你需要用cudaDeviceSynchronize来确保子内核执行完毕。
我的建议：刚开始学动态并行时，别急着写复杂的递归。先跑通最简单的父子内核，确认你的GPU支持动态并行。然后逐步增加复杂度。我曾经见过一个同事，一上来就写八叉树递归，结果三天没调通，最后发现是计算能力不够。
3.6 调试动态并行程序的技巧
动态并行程序的调试比普通CUDA程序难一个数量级。为什么？因为错误发生在GPU内部，你没法打断点。我分享几个经验：
- 用
cudaGetLastError：在父内核里，每次cudaLaunchDevice之后都检查一下返回值。虽然不能精确定位，但至少知道有没有启动失败。 - 限制线程数：调试时，让父内核只用一个线程，子内核也只用一个线程。这样容易跟踪执行流。
- 打印调试信息：在关键位置加
printf，但要注意输出顺序可能乱序。我习惯在每个打印里加上线程ID和层级信息。 - 使用
cuda-memcheck：这个工具能检测动态并行里的内存越界和未初始化变量。强烈推荐。
嗯，说到调试，我想起一个案例。有一次我写了一个动态并行的光线追踪程序，子内核启动子内核，嵌套了5层。结果程序跑起来就卡死。我用cuda-memcheck一查，发现是某个子内核的共享内存越界了。修复后性能提升了30%。所以说，工具用得好，下班下得早。
3.7 本章小结
这一章我们完成了第一个动态并行程序。核心就三点：
- 父子内核概念：父内核从CPU启动，子内核从GPU启动。
- cudaLaunchDevice用法：注意参数传递方式，特别是
args的格式。 - 动态并行的限制：计算能力、嵌套深度、同步问题。
下一章我们会深入动态并行的同步机制，看看如何让父子内核协同工作。说实话，同步才是动态并行最头疼的部分，但也是最精彩的部分。准备好了吗？