必须启用设备端链接（-rdc=true）并链接cudadevrt库，否则cudaLaunchKernel在设备代码中报未定义引用；需指定计算能力≥3.5的架构，子核须用cudaLaunchKernel显式启动，配合__threadfence_system()保证内存可见性，并用cudaDeviceSynchronize()同步。
必须用 nvcc 启用设备端链接（-rdc=true）并显式链接 cudadevrt，否则 cudaLaunchKernel 在设备代码里直接报未定义引用——这是最常卡住的一步。
编译时必须加 -rdc=true 和 -lcudadevrt
动态并行不是默认开启的功能。哪怕 GPU 是 Tesla K80（Compute Capability 3.7）或更新，不加这两个标志，设备端调用 cudaLaunchKernel 或 >> 语法都会失败。
-
-rdc=true启用可重定位设备代码（Relocatable Device Code），让设备函数能跨编译单元被引用 -
-lcudadevrt链接设备端运行时库，提供cudaLaunchKernel、cudaDeviceSynchronize等设备可用符号 - 还必须指定架构，如
-arch=sm_35或更高（Pascal 及以后都支持，但sm_30不行）
完整命令示例：nvcc -arch=sm_50 -rdc=true parent.cu -o parent -lcudadevrt
设备端启动只能用 cudaLaunchKernel，不能用 >> 语法
虽然文档里常写 “父核中可直接用 >> 启动子核”，但实际在 C++ CUDA 项目中，这种写法仅在极简单文件、无分离编译时偶然有效；一旦涉及头文件、多 .cu 文件或类封装，>> 会因符号不可见而链接失败。
立即学习“C++免费学习笔记（深入）”；
- 可靠写法是显式调用
cudaLaunchKernel，传入子核函数指针（强制转为void*）、grid/block 维度、shared memory、stream - 子核函数必须声明在调用前，且不能是模板实例化或 lambda；普通
__global__函数即可 - 注意：设备端没有
cudaGetLastError，错误得靠主机端检查cudaGetErrorString(cudaGetLastError())，或在子核开头加printf排查是否执行
示例片段：
__global__ void child_kernel(float* out) { out[threadIdx.x] = 1.0f; }<br><br>__global__ void parent_kernel(float* d_out) {<br> if (threadIdx.x == 0) {<br> cudaLaunchKernel(<br> (void*)child_kernel,<br> dim3(1), dim3(256),<br> 0, 0, 0<br> );<br> cudaDeviceSynchronize();<br> }<br>}
cudaDeviceSynchronize() 是设备端唯一可用的全局同步原语
设备代码里不能用 cudaStreamSynchronize（参数非法）、也不能用 __syncthreads() 等待子核——后者只同步本块线程，对子核完全无效。
-
cudaDeviceSynchronize()会阻塞当前设备上下文，等待所有已启动的子核（包括嵌套的）完成 - 它代价高，只适合粗粒度控制；频繁调用会严重拖慢性能，尤其当子核本身很轻量时
- 若需异步协作，应改用设备端 stream +
cudaStreamCreateWithFlags（但需确保流对象生命周期在设备端可控，通常只在单次 kernel 中创建/销毁）
子核访问内存时，要注意父核的写操作是否已对子核可见
GPU 不保证设备端 launch 的内存顺序自动刷新。父核写完数据后立即启动子核，子核可能读到旧值。
- 如果父核和子核共享全局内存（比如通过参数传入
float*），必须在 launch 前用__threadfence_system()强制刷出所有写操作 -
cudaDeviceSynchronize()之后的读是安全的，但那是同步点，不是内存屏障 - 避免依赖隐式顺序：不要假设 “launch 前写的内存，子核启动后就一定可见”
典型修复：
__global__ void parent_kernel(float* data) {<br> if (threadIdx.x == 0) {<br> data[0] = 42.0f;<br> __threadfence_system(); // 关键：确保子核能看到这个写<br> cudaLaunchKernel((void*)child_kernel, ...);<br> }<br>}
动态并行真正难的不是语法，而是资源边界：每个子核消耗额外的 GPU 上下文、GMU 调度槽位和设备栈空间。嵌套过深（>3 层）或每秒启动上千个子核，很容易触发 cudaErrorLaunchOutOfResources，而且这种错误不会立刻暴露，可能只在特定负载下才出现。
