cuda dynamic parallel 学习笔记

CUDA动态并行是CUDA编程模型的一个扩展，它允许CUDA内核通过启动新的内核来创建新的线程网格。动态并行性是在开普勒架构中引入的，首次出现在GK110芯片中。在以前的CUDA系统中，内核只能从宿主代码启动。涉及递归、不规则循环结构、时间空间变化或其他不适合平面、单级并行的结构的算法需要通过多个内核发射来实现，这增加了主机的负担和主机设备通信的数量。动态并行支持允许动态发现新工作的算法在不增加主机负担的情况下准备和启动内核。本章描述了CUDA体系结构的扩展功能，它支持动态并行性，包括对CUDA编程模型的修改和添加，以利用这些特性，以及用于开发这种附加容量的指导方针和最佳实践。

背景

许多现实世界的应用程序使用的算法可以动态地改变所执行的工作量。例如，图20.1显示了一个湍流模拟示例，其中所需的建模细节的级别在空间和时间上有所不同。当燃烧流从左向右移动时，活动和强度就会增加。模型右侧模型所需的详细信息级别要比模型左边的要高得多。一方面，使用固定的细网格会导致太多的工作，而在模型的左边没有任何好处。另一方面，使用固定的粗网格在模型的右侧会牺牲太多的准确性。理想情况下，应该为模型的部分使用细网格，这需要更多的细节和粗糙的网格，而不需要太多的细节。

以前的CUDA系统需要从宿主代码中启动所有内核。线程网格所做的工作在内核启动期间是预先确定的。有了内核代码的SPMD编程风格，如果要使用不同的网格间距，就会非常困难。这个限制倾向于使用固定网格系统，正如我们在第12章中讨论的那样。为了达到预期的精度，像图20.1中所示的这样的固定网格方法，通常需要适应模型中最苛刻的部分，并在不需要太多细节的部分执行不必要的额外工作。

图20.1中右下方的动态网格显示了一种更可取的方法。当仿真算法在模型的某些区域检测到快速变化的模拟量时，它对这些区域的网格进行了改进，以达到期望的精度。这样的改进不需要在没有显示出如此密集的活动的领域进行。通过这种方式，算法可以将更多的计算工作动态地引导到从附加工作中受益的模型区域。

图20.2显示了原始CUDA和动态并行版本与图20.1中的模拟模型之间的概念比较。如果没有动态并行性，宿主代码必须启动所有内核。如果发现了新的工作，比如在内核执行期间对模型区域的网格进行细化，那么它需要向主机代码报告，并拥有启动新内核的主机代码。如图20.2（a）所示，主机启动一波内核，接收来自这些内核的信息，并为完成的内核所发现的任何新工作启动下一级内核。

图20.2（b）表明，在动态并行性下，发现新工作的线程可以继续进行，并启动内核来完成工作。在我们的示例中,当一个线程发现需要改进的区域模型,它可以启动一个内核执行计算一步细化网格区域没有终止内核的开销,报告回主机,主机启动新的内核。

【动态launch kernel的好处是，可以减少GPU与CPU之间的通信，】

从程序员的角度来看，动态并行意味着他或她可以在内核中编写内核发射语句。在图20.3中，main函数（主机代码）启动三个内核、A、B和C，这些是原始CUDA模型中的内核发射。不同的是，其中一个内核，B，启动三个内核X，Y，和Z，这在之前的CUDA系统中是非法的。

从内核启动内核的语法与从主机代码中启动内核的语法是一样的：

kernelname<<Dg，Db，Ns，S>（内核参数）

•Dg的类型是dim3，并指定了网格的尺寸和大小。

•Db是一个类型的dim3，并指定每个线程块的大小和大小。

Ns是类型sizet，并指定为这个调用动态分配的共享内存的字节数，这是除了静态分配的共享内存之外的。Ns是一个可选的参数，默认为0。

•S是cudastreamt型的，它指定了与这个调用相关联的流。流必须在发出调用的同一线程块中分配。S是一个可选的参数，默认为0。

20.3重要的细节

尽管从内核发射内核的语法与从主机代码中启动内核的语法类似，但是有几个重要的区别必须被程序员清楚地理解。

发射环境配置

所有的设备配置设置（例如，从cudaDeviceGetCacheConfig()返回的共享内存和L1缓存大小，以及从cudaDeviceGetLimit() 返回的设备限制）将从父kernel继承。也就是说，如果双亲配置为16 K字节的共享内存和48 K字节的L1缓存，那么孩子的执行设置将被配置为相同的。同样地，父类的设备限制，如堆栈大小，将按原样传递给它的子节点。

API错误和启动失败

就像在主机代码中调用CUDA API函数一样，在内核中调用的CUDA API函数可能返回错误代码。最后返回的错误代码被记录下来，并且可以通过cudaGetLastError()调用来检索。错误被记录在每个线程的基础上，这样每个线程都可以识别它所生成的最近的错误。错误代码是cudaError_t 类型，它是一个32位整数值。

Event 事件

只有在内核函数中支持CUDA事件的跨流同步功能。单个流中的事件目前在内核函数中不受支持。这意味着cudaStreamWaitEvent() 是受支持的，但是cudaevent同步()、cudaeventelaptime()的计时，以及通过cudaEventQuery()的事件查询都不是。这些可能会在将来的版本中得到支持。

为了确保这个限制被用户清楚地看到，动态并行的cuda事件必须通过cudaEventCreateWithFlags()来创建，它目前只接受来自内核的cudaEventDisableTiming禁用计时标志值。

事件对象可以在创建它们的CUDA线程块之间共享，但是它们是本地的，并且不应该传递给孩子/父内核。事件句柄不能保证在块之间是唯一的，所以在一个没有分配它的块中使用一个事件句柄将导致未定义的行为。

【shared memory是共享的，但是local memory不共享】

Stream流

【应用程序通过“流”管理上面描述的并发操作。流是按顺序执行的一连串命令（可能是由不同的主机线程发出的）。另一方面，不同的流可以彼此之间或者并发地执行它们的命令;这种行为是没有保证的，因此不应该依赖于正确性（例如，内核间通信是未定义的）。】

在动态并行性下，命名和未命名（NULL）流都是可用的。命名流可以被线程块中的任何线程使用，但是流句柄不应该传递给其他块或子/母内核。换句话说，一个流应该被看作是私有的，在它被创建的块中。流句柄不能保证在块之间是唯一的，所以在一个没有分配它的块中使用流句柄将导致未定义的行为。

与主机端启动类似，在单独的流中启动的工作可以并发运行，但是实际的并发性并不能保证。在子内核之间需要并发的程序是不完善的，并且会有未定义的行为。

在动态并行性下，主机端空流的全局同步语义不受支持。为了显式地指出这种行为变化，必须使用cudaStreamCreateWithFlags（）API在内核中使用cudastreamnon阻塞标志来创建所有流。对cudaStreamCreate（）的调用将会由于编译器“未被识别的函数调用”错误而失败，以便在动态并行性下明确不同的流语义。

cudastream同步（）API在内核中是不可用的;只有cudadevicesyn同步（）可以被用来显式地等待启动的工作完成。这是因为底层系统软件只实现了块范围的同步调用，并且不希望提供一个不完整语义的API（例如同步保证了一个流同步，但碰巧提供了一个完整的屏障作为副作用）。

一个执行网格的一部分，配置和启动一个新网格的线程属于父网格，而由发射创建的网格是子网格。如图20.4所示，子网格的创建和完成是正确嵌套的，这意味着在所有由线程创建的子网格完成之前，父网格不被认为是完整的。即使父线程在启动的子网格上没有显式地同步，运行时保证父节点和子节点之间的隐式同步，强制父进程在退出执行之前等待它的所有子节点退出执行。

参考文献：Programming Massively Parallel Processors A Hands-on Approach