同步嵌套内核的深度

Synchronizing depth of nested kernels

本文关键字:深度 内核 嵌套 同步      更新时间:2023-10-16

让我们在有父内核和子内核的地方采用以下代码。从上述父内核开始,我们希望在不同的流中启动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();

在父内核完成之前,父内核将等待任何生成的子内核完成。 动态并行性文档中对此进行了介绍:

子网格

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

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

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

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

在 CDP 环境中需要注意的其他几个限制是嵌套限制和挂起启动限制。