同步嵌套内核的深度

Question

让我们在有父内核和子内核的情况下使用以下代码。从所述父内核开始，我们希望在不同的流中启动threadIdx.x 子内核，以最大化并行吞吐量。然后我们等待那些带有cudaDeviceSynchronize() 的子内核，因为父内核需要查看对global 内存所做的更改。

现在假设我们还希望使用流启动 n 父内核，并且在我们希望并行启动的每组 n 父内核之间，我们还必须使用 cudaDeviceSynchronize() 等待结果

这将如何表现？

来自this official introduction to Dynamic Parallelism by 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();
}
for (int i=0; i<10; i++)
{
    cudaStream_t s;
    cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);

    parent_kernel <<<10,10,0,s>>> ();
    cudaStreamDestroy(s);
}
cudaDeviceSynchronize();

Answer 1

在父内核完成之前，父内核将等待任何生成的子内核完成。这在dynamic parallelism documentation:

子网格的调用和完成是正确嵌套的，这意味着在其线程创建的所有子网格都完成之前，父网格不会被认为是完整的。即使调用线程没有在启动的子网格上显式同步，运行时也会保证父子网格之间的隐式同步。

任何其他语义都应该可以从普通流语义中推断出来，即：启动到特定流中的活动将不会开始，直到所有先前启动到该流中的活动都完成。同样，启动到单独流中的活动之间也没有强制排序。

在您的示例中（或实际上在任何示例中），父内核将等待，直到从该父内核启动的所有子内核都完成，无论使用或未使用哪些流。

不清楚您是在问这个问题，但请注意，对于您示例中的设备代码，cudaDeviceSynchronize() 仅保证 该线程 将等待子内核完成，同样仅为该线程强制执行结果可见性排序。如果您希望同一块中的其他线程能够见证线程 0 生成的子内核的全局内存结果（仅举一个例子），那么您需要在线程 0 中使用 cudaDeviceSynchronize() 操作来跟进__syncthreads() 操作。在__syncthreads() 之后，同一块中的其他线程将保证对线程 0 启动的子内核（或任何线程启动的子内核，然后是 cudaDeviceSynchronize() 调用，之前的调用）产生的全局结果的可见性前面提到的__syncthreads())。

在 CDP 环境中需要注意的其他几个限制是 nesting limit 和 pending launch limit。