如何在CUDA父内核中安全同步子内核，确保其完成后再执行后续计算？

首先，你的代码存在基础错误：__device__修饰的函数是设备端普通函数，不能用<<<>>>内核启动语法调用。如果要把child作为子内核启动，必须将其声明为__global__；如果只是需要设备端函数调用，直接调用即可，无需启动——这种情况下函数执行是同步的，父线程会等待child执行完成后再进入后续代码，不需要额外同步。

下面分两种场景给出正确实现：

场景1：child作为设备端普通函数（无需独立内核） #

如果不需要child作为独立内核运行，直接调用设备函数即可，天然满足同步要求：

__device__ void child(double* A, double* B, int p)
{
    int r = blockIdx.x * blockDim.x + threadIdx.x;
    B[p + r] += A[r];
}

__global__ void parent(double* A, double* B, double* C)
{
    int i, r = blockIdx.x * blockDim.x + threadIdx.x;
    // 直接调用设备函数，执行完成后才会继续后续代码
    child(A, B, 7 * r);
    for (i = 0; i < 5; i++)
        C[r] += B[r * 7 + i];
}

int main()
{
    parent<<<1, 2>>>(A, B, C);
}

场景2：child必须作为独立子内核（使用动态并行） #

如果确实需要child作为独立内核运行（比如需要更大的线程网格），则需要启用CUDA动态并行特性，同时在父内核中使用设备端同步原语等待子内核完成：

前置要求： #

• 使用计算能力3.5及以上的GPU（支持动态并行）
• 编译时添加参数-rdc=true（启用可重定位设备代码）和-lcudadevrt（链接CUDA设备运行时库）

基础同步实现代码： #

// 将child声明为__global__，作为子内核
__global__ void child(double* A, double* B, int p)
{
    int r = blockIdx.x * blockDim.x + threadIdx.x;
    B[p + r] += A[r];
}

__global__ void parent(double* A, double* B, double* C)
{
    int i, r = blockIdx.x * blockDim.x + threadIdx.x;
    // 启动子内核
    child<<<1, 5>>>(A, B, 7 * r);
    // 设备端同步：等待所有子内核完成
    cudaDeviceSynchronize();
    // 子内核执行完毕后，再进行后续计算
    for (i = 0; i < 5; i++)
        C[r] += B[r * 7 + i];
}

int main()
{
    parent<<<1, 2>>>(A, B, C);
    // 主机端等待父内核完成
    cudaDeviceSynchronize();
}

更高效的流同步方式（可选） #

如果不想同步整个设备，可使用设备端流限制同步范围，减少性能损耗：

__global__ void child(double* A, double* B, int p)
{
    int r = blockIdx.x * blockDim.x + threadIdx.x;
    B[p + r] += A[r];
}

__global__ void parent(double* A, double* B, double* C)
{
    int i, r = blockIdx.x * blockDim.x + threadIdx.x;
    cudaStream_t stream;
    // 在设备端创建私有流
    cudaStreamCreate(&stream);
    // 将子内核启动到指定流中
    child<<<1, 5, 0, stream>>>(A, B, 7 * r);
    // 仅同步该流，等待流内的子内核完成
    cudaStreamSynchronize(stream);
    // 销毁流
    cudaStreamDestroy(stream);
    
    for (i = 0; i < 5; i++)
        C[r] += B[r * 7 + i];
}

int main()
{
    parent<<<1, 2>>>(A, B, C);
    cudaDeviceSynchronize();
}

内容的提问来源于stack exchange，提问作者AbdelAziz AbdelLatef