同步嵌套内核的深度
Synchronizing depth of nested kernels
让我们在有父内核和子内核的地方采用以下代码。从上述父内核开始,我们希望在不同的流中启动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 环境中需要注意的其他几个限制是嵌套限制和挂起启动限制。
- C++17复制构造函数,在std::unordereded_map上进行深度复制
- 如何在内核C++中使用1920x1080x16M图形或类似的16M颜色?(VGA)
- OpenGL在启用深度测试时不会丢弃我的碎片
- CUDA内核和数学函数的显式命名空间
- 码头化的C++应用程序是否向后兼容早期的内核版本
- C++内核出现Jupyter笔记本错误
- 当我尝试加载内核模块时,如何修复C++中的这个 malloc() 错误?
- C++尝试深度复制唯一指针时出现内存访问冲突
- 内存围栏是否涉及内核
- 将 2D 推力::d evice_vector 复矩阵传递给 CUDA 内核函数
- 试图找到二叉树的深度
- OpenCL 内核参数中的字符***?
- 具有可分离内核的 2D 模糊卷积
- 如何在Windows内核中获取文件大小
- 库达如何将字符**从内核复制到主机
- 操纵安卓相机的深度图导致应用程序崩溃
- 深度值没有意义 R200 相机
- OpenCL 是否支持向量作为内核参数?
- 对如何制作双链表的深度副本感到困惑?
- 同步嵌套内核的深度