CUDA动态并行中如何让父核等待子核完成后再执行后续代码？

在CUDA动态并行场景下，需要子核runParatron完全执行完毕后，父核runMLP才进入for循环的下一次迭代。但当前输出显示子核的打印语句晚于父核的“[”打印，说明未实现等待。尝试调用cudaDeviceSynchronize函数时编译报错，提示该主机代码无法在设备代码中执行且未定义，求正确的同步方法。

相关代码：

__global__ void runMLP(double* x, double* outputs, double* weights, activation_function* A_Fs, int*     CIL, int layers, int bias, int* WLO, int* OLO) {

    if (CIL[0] > 511) {
        copyElements <<<CIL[0] / 32, 32>>> (outputs, x, CIL[0]);
        //I WOULD ALSO LIKE TO WAIT HERE
    }
    else
        for (int i = 0;i < CIL[0];i++) {
            outputs[i] = x[i];
        }

    for (int i = 1;i < layers;i++) {
        printf("----------------------Layer %d :: InputSize %d :: Layer weight offset %d ::     Layer output offset %d----------------------\n", i, CIL[i-1], WLO[i-1], OLO[i]);
        runParatron <<< (CIL[i] / 32) + 1, 32>>> (outputs + OLO[i - 1], outputs +     OLO[i], weights + WLO[i - 1], A_Fs[i], CIL[i - 1], CIL[i], bias);
        //cudaDeviceSynchronize(); //THIS IS WHERE I NEED TO WAIT UNTIL NEXT ITERATION
    }
    if (A_Fs[layers - 1] == SOFTMAX) {
        double* temp = outputs + OLO[layers - 1];
        printf("[");
        for (int i = 0;i < CIL[layers-1];i++) {
            printf("% d, ", temp[i]);
        }
        printf("]\n");
        double denom = 0;
        for (int i = 0;i < CIL[layers - 1];i++) {
            denom += temp[i];
        }
        if (denom < DBL_MIN)
            denom = DBL_MIN;
        for (int i = 0;i < CIL[layers - 1];i++) {
            temp[i] /= denom;
        }
    }
}

异常输出示例：

//All Cell: starting lines are produced from child kernel
[Cell: 0 :: weightOffset 0 :: AF 2 //As you can see, there is the "[" here when it should be
Cell: 1 :: weightOffset 6 :: AF 2
Cell: 2 :: weightOffset 12 :: AF 2
Cell: 3 :: weightOffset 18 :: AF 2
-502657059,  2118981138,  1645236453, ] //Down here!

• 正确启用cudaDeviceSynchronize()：设备端调用该函数需要两个前提：编译时添加-rdc=true选项启用动态并行，并链接CUDA runtime库；GPU计算能力需达到3.5及以上（支持动态并行）。之前编译报错大概率是未开启动态并行编译选项。
• 流同步替代方案：可以在父核内创建CUDA流，将子核任务提交到该流后，调用cudaStreamSynchronize等待任务完成。示例代码如下：

  for (int i = 1;i < layers;i++) {
      printf("----------------------Layer %d :: InputSize %d :: Layer weight offset %d ::     Layer output offset %d----------------------\n", i, CIL[i-1], WLO[i-1], OLO[i]);
      cudaStream_t stream;
      cudaStreamCreate(&stream);
      runParatron <<< (CIL[i] / 32) + 1, 32, 0, stream>>> (outputs + OLO[i - 1], outputs + OLO[i], weights + WLO[i - 1], A_Fs[i], CIL[i - 1], CIL[i], bias);
      cudaStreamSynchronize(stream);
      cudaStreamDestroy(stream);
  }
  

• 性能与兼容性注意：动态并行同步会带来一定延迟，需评估对整体性能的影响；同时要确认运行代码的GPU支持动态并行（计算能力3.5+）。

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