Cuda动态并行性:可以创建的子线程深度

问题描述

我正在阅读CUDA编程指南,我发现它很密集。我在本节中向他们解释了父网格可以创建子网格,并且仅当其所有生成的子线程都已完成时,才认为父网格已完成。

我的问题是:允许在Cuda中生长的父子树有多“深”:这些仅受所讨论硬件的计算能力的约束,例如,可以生成尽可能多的父/子块吗?只要我们不超过可以一次在硬件上运行的最大线程数,还是他/她想要的最大线程数,还是有进一步的限制?我之所以这样问是因为缺少该功能,我看不到如何在GPU上实现递归。

谢谢, 胺

解决方法

我的问题是:允许在Cuda中生长的父子树有多“深”

documentation表示最大嵌套深度为24。

如文档中所示,通常在实际达到24的嵌套深度之前,您可能会首先遇到其他限制。其中之一是设备内核启动的一般限制,包括内存要求以及待决的启动限制。另一个可能的限制是同步限制。这与父内核是否显式等待子内核完成(例如通过设备端cudaDeviceSynchronize())以及此同步扩展到什么深度有关。

只要我们不超过可以一次在硬件上运行的最大线程数

这都不明显取决于父内核或子内核中有多少个线程。 CUDA内核对硬件可以一次运行的线程数没有任何基本限制,CUDA动态并行性(CDP)也没有。

那么,实际上,大深度CDP发射可能遇到各种限制。此外,从性能的角度来看,这样的设计模式可能不是最好的。 CDP启动具有与之相关的时间和资源开销,对于以这种方式细分工作的任何模式,在CUDA内核中通常希望在内核中进行更多的工作,而不是更少。 / p>