嵌套CUDA内核的同步深度与执行行为问询

场景与问题 #

我遇到了这样一个CUDA编程场景，想请教下相关行为：

• 父内核中会启动threadIdx.x个子内核到不同流中，以此最大化并行吞吐量，之后调用cudaDeviceSynchronize()等待子内核完成，因为父内核需要感知全局内存的修改。
• 现在我在不同流中启动n个父内核，且每轮并行启动的n个父内核之间需要调用同步等待结果，想知道程序会有怎样的表现？
• 根据NVIDIA官方动态并行文档，我推测parent kernel[0]仅会等待其内部启动的流，这个推测是否正确？如果不正确，实际行为是什么？

注：已知同时运行的流数量有限（本文设定为32个），这么做的目的是最大化设备占用率。

代码示例 #

__global__ void child_kernel (void) {}

__global__ void parent_kernel (void) {
    if (blockIdx.x == 0) {
        cudaStream_t s;
        cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
        child_kernel <<<1,10,0,s>>> ();
        cudaStreamDestroy(s);
    }
    cudaDeviceSynchronize();
}

int main() {
    for (int i=0; i<10; i++) {
        cudaStream_t s;
        cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
        parent_kernel <<<10,10,0,s>>> ();
        cudaStreamDestroy(s);
    }
    cudaDeviceSynchronize();
}

详细解答 #

先给你明确结论：你的推测不正确，实际行为和你预想的有很大差异——父内核里的cudaDeviceSynchronize()会等待整个GPU设备上所有未完成的任务，而不仅仅是它自己启动的子流。

具体拆解下原因和影响：

1. 动态并行中同步函数的范围：在CUDA动态并行机制里，内核内部调用的cudaDeviceSynchronize()是全局级别的同步操作。它会阻塞当前父内核的执行，直到整个GPU上所有正在运行的工作（包括其他父内核、其他父内核启动的子内核，甚至主机端其他流中提交的任务）全部完成。
2. 你的代码场景中的问题：
   • 你在主机端循环将10个父内核提交到不同的非阻塞流中，原本是希望这些父内核并行执行来最大化设备占用率，但每个父内核里的全局同步直接打破了这个计划。
   • 只要有一个父内核执行到cudaDeviceSynchronize()，它就会等待所有其他9个父内核以及它们的子内核全部完成才能继续；反过来，其他父内核执行到这个同步点时也会做同样的操作。最终的结果就是这些父内核根本无法并行，只能串行执行，完全达不到你想要的提升占用率的目标。
3. 正确的同步方式：如果想实现你最初的目标（父内核仅等待自己启动的子内核，不影响其他父内核的并行执行），你应该把父内核里的cudaDeviceSynchronize()替换为cudaStreamSynchronize(s)，只同步你在父内核内部创建的那个子流。这样每个父内核只会等待自己启动的子内核完成，其他流中的父内核可以正常并行运行，才能真正利用并行性提升设备占用率。
4. 流数量限制的影响：你提到同时运行的流数量上限为32个，只要主机端创建的流不超过这个限制，原本是可以支撑父内核并行的，但全局同步直接废掉了这个并行优势。

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