同步嵌套内核的深度

Question

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

现在假设我们还希望使用流启动 n parent 内核，并且在我们希望并行启动的每组 n parent 内核之间，我们还必须使用 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:

The invocation and completion of child grids is properly nested, meaning that the parent grid is not considered complete until all child grids created by its threads have completed. Even if the invoking threads do not explicitly synchronize on the child grids launched, the runtime guarantees an implicit synchronization between the parent and child.

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

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

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

之前

在 CDP 环境中需要注意的其他一些限制是 nesting limit and the pending launch limit。

同步嵌套内核的深度

Synchronizing depth of nested kernels

c++

cuda

dynamic-parallelism