0

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

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

这将如何表现？

从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();
4

1 回答 1

3

在父内核完成之前，父内核将等待任何生成的子内核完成。这在动态并行文档中有所介绍：

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

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

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

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

在 CDP 环境中需要注意的其他一些限制是嵌套限制和待定启动限制。

于 2018-05-22T23:57:54.447 回答