环境与编译：支持动态并行的GPU架构（Kepler+），编译选项（-rdc=true），运行时API

好，咱们正式开始动手之前，先把环境搭好。这块儿其实挺关键的，很多人代码写好了，一编译报错，查半天发现是编译选项没开对。我当年第一次搞动态并行就栽在这上面，折腾了整整一个下午。

硬件要求：Kepler架构及以后

动态并行这个特性，不是所有GPU都支持的。它最早出现在NVIDIA的Kepler架构上，也就是Compute Capability 3.5及以上的设备。

说白了，你得有一块比较新的显卡。具体来说：

• Kepler架构（CC 3.5）：GTX 780、Tesla K40 等
• Maxwell架构（CC 5.x）：GTX 980、Tesla M40 等
• Pascal架构（CC 6.x）：GTX 1080、Tesla P100 等
• Volta架构（CC 7.0）：Tesla V100 等
• Turing架构（CC 7.5）：RTX 2080 等
• Ampere架构（CC 8.x）：RTX 3080、A100 等

怎么查你的显卡支不支持？很简单，跑一下 deviceQuery 示例程序，看输出里的 Compute Capability 字段。我习惯把这个程序放在手边，每次换机器都先跑一遍确认。

编译选项：-rdc=true 是必须的

这是个大坑，也是很多人第一次搞动态并行时最容易忽略的地方。

正常情况下，你编译CUDA代码用 nvcc，默认是生成一个单独的设备代码对象。但动态并行需要在设备端也能调用kernel，这就涉及到设备代码的链接问题。

所以，必须加上 -rdc=true（或者 --relocatable-device-code=true）。这个选项告诉编译器：生成可重定位的设备代码，允许设备端函数被其他设备函数调用。

来看一个典型的编译命令：

nvcc -arch=sm_35 -rdc=true -o my_program my_program.cu

嗯，这里要注意几点：

• -arch=sm_35 指定目标架构，至少是3.5
• -rdc=true 必须加，否则链接时会报错
• 如果你用了多个.cu文件，还需要链接 cudadevrt 库

多文件编译时，命令会变成这样：

nvcc -arch=sm_35 -rdc=true -c kernel1.cu -o kernel1.o
nvcc -arch=sm_35 -rdc=true -c kernel2.cu -o kernel2.o
nvcc -arch=sm_35 -rdc=true -o my_program kernel1.o kernel2.o -lcudadevrt

看到最后那个 -lcudadevrt 了吗？这是设备运行时库，动态并行全靠它支撑。我刚开始做项目时，忘了加这个，链接报了一堆未定义符号的错误，当时差点以为编译器坏了。

运行时API：cudaSetDevice 与流管理

动态并行在运行时层面也有一些特殊要求。最核心的一点是：设备端启动的kernel，默认使用与父kernel相同的流和设备上下文。

但有时候，你想在设备端指定不同的流，或者在不同的设备上启动子kernel。这时候就需要用到运行时API了。

举个例子，在设备端获取当前设备ID：

__global__ void parent_kernel() {
    int device_id;
    cudaGetDevice(&device_id);
    printf("Parent kernel running on device %d\n", device_id);
    
    // 启动子kernel
    child_kernel<<<1, 256>>>();
}

你想想看，这个 cudaGetDevice 在设备端也能用，是不是挺方便的？但要注意，不是所有运行时API都能在设备端调用。只有那些标记为 设备端可调用 的API才行。

常用的设备端运行时API包括：

API函数	功能说明	设备端可用
cudaGetDevice	获取当前设备ID	是
cudaSetDevice	设置当前设备	是
cudaStreamCreate	创建流	是
cudaStreamSynchronize	同步流	是
cudaMalloc	分配设备内存	是
cudaFree	释放设备内存	是
cudaMemcpy	内存拷贝	否
cudaEventCreate	创建事件	是

看到没？cudaMemcpy 在设备端是不能用的。为什么？因为设备端不能直接访问主机内存，这个操作在硬件层面就不支持。我曾经试图在子kernel里用 cudaMemcpy 把数据拷回主机，结果编译通过了，运行时直接挂掉。嗯，从那以后我就记住了这个限制。

避坑指南：我踩过的几个坑

最后，分享几个我实际项目中遇到的坑，希望能帮你省点时间：

• 坑一：忘记加 -rdc=true —— 编译报错 "error: calling a __global__ function from a __global__ function is only allowed with --relocatable-device-code=true"。这个错误信息其实挺明确的，但新手容易忽略。
• 坑二：动态并行深度限制 —— Kepler架构最多支持24层嵌套，Pascal以后放宽到64层。但别真去试64层，我试过，性能惨不忍睹。
• 坑三：设备端内存泄漏 —— 在设备端用 cudaMalloc 分配的内存，必须在设备端用 cudaFree 释放。如果父kernel结束了还没释放，这块内存就永远找不回来了。
• 坑四：同步问题 —— 子kernel默认是异步启动的，父kernel不会等子kernel执行完就继续往下走。如果需要等待子kernel完成，记得加 cudaDeviceSynchronize。

下一章，我们会真正动手写第一个动态并行的例子。到时候你会发现，其实代码写起来并不复杂，关键是理解背后的执行模型。准备好了吗？