7.5 动态并行

动态并行（Dynamic parallelism）是一个仅适用于SM3.5架构硬件的新功能，可以使CUDA内核启动其他CUDA内核，同时也可以在CUDA运行时中调用各种函数。当使用动态并行时，CUDA运行时的一个子集（即所谓的“设备运行时”）可以为设备上运行的线程使用。

动态并行引入“父”和“子”网格的术语。任何被另一个CUDA内核启动的内核（而不是主机代码，正如所有以前的CUDA版本）称为一个“子内核”，而调用它的为“父内核”。默认情况下，CUDA支持两级嵌套（一个给父网格，另一个给子网格），这个数目可以通过调用带cudaLimitDevRuntimeSyncDepth标志的CUDASetDeviceLimit()函数来增加。

动态并行为处理应用程序而设计。而在它出现之前GPU必须将结果发送给CPU以便于CPU可以指定在GPU上要执行哪些工作。这样的“握手”会扰乱第2.5.1小节中描述的执行流水线中的CPU/GPU并发，其中CPU会产生为GPU使用的命令。GPU的时间很宝贵，GPU不会在分配到更多工作之前一直等待CPU读取和分析结果。动态并行会让GPU在内核中为自己启动任务以避免这些流水线中的“气泡”。

动态并行可以提高如下性能。

·在内核可以开始执行之前，给内核所需要的数据结构进行初始化。在此之前，这样的初始化要在主机代码中处理，或由之前启动的一个单独内核处理。
·在诸如Barnes-Hut引力积分或空气动力学模拟的分层网格评估等应用中，它可以简化递归。

注意 动态并行只在一个给定GPU中有效。内核可以调用内存复制操作或其他内核，但却无法把工作提交给其他GPU。

7.5.1 作用域和同步

子网格会继承其父网格的线程块和网格大小之外的绝大部分内核配置参数，如共享内存配置（由CUDADeviceSetCacheConfig()设置）。线程块是作用域的基本单位：一个线程块创建的流和事件仅能被该线程块使用（它们甚至不会被子网格继承），并且它们会在该线程块退出时自动销毁。

注意 通过动态并行在设备上创建的资源与在主机上创建的资源是严格分开的。在主机上创建的流和事件不能通过动态并行在设备上使用，反之亦然。

CUDA确保在所有子网格完成之前，父网格是不会完成的。虽然父网格与子网格可以并发执行，但这却不能保证一个子网格会在它的父网格调用CUDADeviceSynchronize()之前就开始执行。

如果线程块中的所有线程都退出，该线程块的执行将会被挂起，直到所有子网格完成为止。如果默认同步不够用，开发者可以使用CUDA流和事件来显式同步。与主机上一样，在一个给定流中的操作以子任务的顺序执行。当这些操作被指定到不同的流中，它们就只能并发执行了，并且还不能保证它们实际上会并发执行。如果需要，可以使用__syncthreads()等同步原语来协调给定流中子任务的顺序。

注意 在设备上创建的流和事件不能在创建它们的线程块之外使用。

CUDADeviceSynchronize()函数对由块中任何线程启动的挂起工作进行同步。然而，它不执行任何线程间的同步。因此，如果要在其他线程启动的工作上同步的话，开发者必须使用__syncthreads()或其他块级同步原语（详见8.6.2小节）。

7.5.2 内存模型

父与子网格共享相同的全局与常量内存存储，但是它们有不同的本地和共享内存。

1. 全局内存

如下两种情况下，子网格在执行时，其内存视图与父网格具有完全的一致性：子网格被父网格调用，以及子网格被父线程中同步API调用产生的信号导致结束。

所有父线程中的全局内存操作对子网格来说是可见的。父网格与子网格同步完成后，所有子网格中的内存操作对其父网格来说是可见的。零复制内存具有与全局内存相同的连贯性和一致性保证。

2. 常量内存

常量是不变的，且不得在内核执行时从设备中进行修改。从一个内核线程内获取一个常量内存对象的地址与所有的CUDA程序具有一样的语义 [1]，并完全支持在父网格与子网格之间传递该指针。

3. 共享与本地内存

共享和本地内存对于线程块和线程分别是私有的，并且在父与子网格之间不是可见或一致的。当这些位置的一个对象于它的作用域之外被引用，该行为是未定义的并且有可能会导致错误。

如果nvcc检测到一个滥用共享或本地内存指针的操作，它将会发出警告。开发者可以使用内建__isGlobal()来确定一个给定指针是否被全局内存引用。共享或本地内存的指针对于CUDAMemcpyAsync()或CUDAMemsetAsync()来说是无效的参数。

4. 本地内存

本地内存是一个处于执行中的线程的私有存储空间，且在该线程之外是不可见的。启动子内核时传递一个指向本地内存的指针作为启动参数是非法的。从一个子内核中解引用这样一个本地内存指针的结果将会是未定义的。要保证这个规则不会在无意中被编译器引用，所有传递给子内核的存储空间应从全局内存堆上进行显示分配。

5. 纹理内存

父和子内核的并发访问有可能会导致数据的不一致，因此，应当避免这种行为。也就是说，父与子内核之间的一致程度由运行时决定。一个子内核可以使用纹理操作来访问其父内核写入的内存，但子内核的内存写操作并不会在父内核的纹理内存访问中反映出来，直到父与子完成同步之后为止。纹理对象在设备运行时上得到很好支持。它们不能够被创建或销毁，但可以进行传递和在层次结构中被任何网格使用（例如父和子网格）。

7.5.3 流与事件

设备运行时所创建的流和事件仅能在创建它们的线程块中使用。NULL流在设备运行时与在主机运行时具不同语义。在主机上，与NULL流的同步会强制让GPU上其他所有的数据流操作一同“汇聚”（如第6.2.3节中描述的那样）；而在设备上，NULL流是一条独立的流，并且任何流之间的同步都必须使用事件进行。

在使用设备运行时的时候，流必须以cusadStreamNonBlocking标志（cusadStreamCreateWithFlags()的一个参数）创建。

CUDAStreamSynchronize()调用是不支持的，要同步的话必须根据事件和CUDAStreamWaitEvent()进行。

CUDA事件中，只有流之间的同步功能是支持的。其结果是 CUDAEventSynchronize()、CUDAEventElapsedTime()和 CUDAEventQuery()不被支持。此外，由于不支持计时功能，事件必须通过将CUDAEventDisableTiming标志传递给 CUDAEventCreateWithFlags()来创建。

7.5.4 错误处理

设备运行时上的任何函数都有可能会返回一个错误（CUDAError_t）。该错误会被记录在线程私有的某插槽中，并且可以

通过调用CUDAGetLastError()来进行查询。与基于主机的运行时一样，CUDA会区分可以立即返回的错误（例如，如果一个无效参数被传递到一个内存复制函数）和必须以异步方式返回的错误（例如，如果启动执行了一个无效的内存访问）之间的不同。如果一个子网格在运行时导致一个错误，CUDA会返回一个错误给主机而不是父网格。

7.5.5 编译和链接

与主机运行时不同，在使用设备运行时的时候，开发者必须对设备运行时的静态库进行显式链接。在Windows中，设备运行时为cudadevrt.lib；而在Linux和MacOS上，为cudadevrt.a。在用nvcc构建时，这可能会通过给命令行追加lcudadevrt来完成。

7.5.6 资源管理

无论一个内核什么时候启动子网格，该子网格都将会被视为一个新的嵌套层次，而层次的总数视为该程序的嵌套深度。与此相反，子网格启动时程序显式同步的最深层次被称为“同步深度”。通常情况下，同步深度比程序的嵌套深度小1，但如果程序并不总是需要调用CUDADeviceSynchronize()的话，那么它可能会远小于嵌套深度。

理论上的最大嵌套深度为24，但实际上，它受设备限制项CUDALimitDevRuntimeSyncDepth所控制。深于最大深度的任何内核启动都将失败。默认的最大同步深度是2。而该限制必须在最顶层的内核从主机上启动之前配置好。

注意 调用一个形如CUDAMemcpyAsync的设备运行时函数可以启动一个内核，同时嵌套深度增加1。

对于从不调用CUDADeviceSynchronize()的父内核，系统并不需要为父内核预留空间。在这种情况下，一个程序所需的内存使用量将远远小于预留的最大值。这种程序可以指定一个较小的最大同步深度，以避免过度分配后备存储。

1. 内存使用量

设备运行时的系统软件会预留设备内存以作如下用途。

·跟踪挂起网格的启动；
·在同步过程中保存父网格状态;

• 为内核调用的malloc()和CUDAAlloc()充当可分配的堆。

这一内存不能被应用程序使用，所以有些应用程序可能会降低默认分配，还有某些应用程序可能会为了操作准确而增加默认值。要更改默认值，如表7-3中总结的那样，开发者可以调用CUDADeviceSetLimit()。限制项CUDALimitDevRuntimeSyncDepth是特别重要的，因为每一个嵌套层次将消耗高达150MB的设备内存。

2. 挂起的内核启动

在内核启动时，所有相关的配置和参数数据都会被跟踪，直到内核完成为止。这批数据被存在一个系统管理的启动池中。该池的大小是可通过从主机中调用CUDADeviceSetLimit()，并指定CUDALimitDevRuntimePendingLaunchCount来配置。

3. 配置选项

设备运行时系统软件的资源分配是通过从主机程序的CUDADeviceSetLimit()API来控制的。配置必须在任何内核启动之前设置，并且在GPU运行中不再改变。

设备运行时分配的内存必须由设备运行时释放。此外，由设备运行时分配的内存来自预分配堆，其大小由设备限制项

CUDALimitMallocHeapSize指定。表7-3中列出的限制项可根据需要设置。

表7-3 CUDADeviceSetLimit()值表

7.5.7 小结

表7-4总结了设备运行时和主机运行时之间的主要差异和限制。表7-5中列出了有可能被设备运行时调用的功能子集，以及所有相关限制。

表7-4 设备运行时限制

表7-5 CUDA设备运行时函数

[1] 请注意，在设备代码中，地址必须采用“地址”运算符（一元运算符&），这是因为CUDAGetSymbolAddress()并不支持设备运行时。