CUDA中如何在一个核函数内等待另一个核函数执行完成？

问题本质 #

你遇到的是CUDA动态并行的异步执行问题：父核启动子核后会立即继续执行后续代码，不会默认等待子核完成。__syncthreads()仅用于同步同一个线程块内的线程，完全无法解决跨核函数的同步需求。

正确解决步骤 #

1. 添加同步调用 #

在父核中启动子核的代码之后，调用cudaDeviceSynchronize()，该函数会阻塞当前核函数的线程，直到设备上所有未完成的操作（包括子核的执行）全部结束：

__global__ void test_cmp()
{
    int* a;
    cudaMalloc(&a, sizeof(int));
    *a = 2;
    int* b;
    cudaMalloc(&b, sizeof(int));
    *b = 3;
    int* cmp;
    cudaMalloc(&cmp, sizeof(int));
    *cmp = -3;

    printf("test_cmp cmp %d\n", *cmp);
    cmp_int<<<1, 1>>>(a, b, cmp);
    // 等待子核执行完成
    cudaDeviceSynchronize();
    printf("test_cmp cmp %d\n", *cmp);

    cudaFree(a);
    cudaFree(b);
    cudaFree(cmp);
}

2. 编译配置 #

启用动态并行需要特定的编译选项（以nvcc为例）：

nvcc your_code.cu -o output -rdc=true -lcudadevrt

• -rdc=true：开启可重定位设备代码编译，支持核函数内的动态并行调用
• -lcudadevrt：链接CUDA设备运行时库，提供动态并行的底层支持

3. 硬件要求 #

你的GPU必须支持计算能力3.5及以上（Kepler架构及后续型号），这是CUDA动态并行的最低硬件要求。

额外提示 #

• 核函数内的cudaMalloc/cudaFree属于动态内存操作，建议添加错误检查（比如判断cudaMalloc的返回值是否为cudaSuccess），避免内存分配失败引发的隐性问题。
• 动态并行会引入额外性能开销，如果只是简单的数值比较逻辑，直接在父核内实现会比启动子核更高效。

修改并正确编译运行后，父核的第二个printf会拿到子核修改后的cmp值，输出符合你的预期（printf的顺序可能因设备调度略有变化，但核心逻辑的结果是正确的）。

内容的提问来源于stack exchange，提问作者P'tit Ju