CUDA10.0官方文档的翻译与学习之CUDA动态并行

因为启动配置可以在运行于设备上的线程所决定，那么这种直接从GPU上创建新任务就减少了主机和设备间的执行控制和数据传输。另外，数据依赖的并行任何可以在运行时和核函数内联生成，动态利用GPU的硬件调度器和负载均衡器，适应了数据驱动的决策或工作。之前为消除递归、不规则循环结构，或者没有适应扁平化的单层并行度而要求修改的算法和编程模型可以表达得更加透明，但是动态并行只支持计算能力>=3.5的设备

术语

定义一些本章要用到的一些术语，见下表

执行环境和内存模型

执行环境

CUDA执行环境是基于线程、线程块、网格和定义被线程块和网格中每个线程执行程序的核函数原语的，当一个核函数启动时，网格的属性就被在CUDA中有特殊语法的执行配置描述了。CUDA中对动态并行的支持为在设备上已经运行的线程扩展了配置、启动和同步新网格的能力。

父子网格

配置启动新网格的网格属于父网格，被创建的叫做子网格。子网格的创建和完成是无缝嵌套的，这意味着要等到其内线程创建的所有子网格都完成了，父网格才能算完成。即便调用线程没有显式在子网格启动时调用同步，运行时也会在父子网格之间隐式地保持同步，如下图所示

CUDA原语的范围

在主机端和设备端，CUDA运行时提供了启动核函数、等待已启动的核函数完成以及通过流和事件追踪启动间的依赖的API。在主机端，进程内所有线程都能访问启动状态和涉及流与事件的CUDA原语；然而，进程时独立执行的，不能共享CUDA对象。同样的层次也存在于设备端，已启动的核函数和CUDA对象对线程块内的所有线程是可见的，但是在线程块之间则是不可见的。这意味一个流可以被一个线程创建，并被同一个块中另一个线程使用，但是不允许和在别的线程块中的线程共享。

同步

任何线程的CUDA运行时操作，包括核函数启动都是在线程块内可见的，这意味着在父网格中的执行线程可以在它自己创建的、线程块中别的线程创建的或者线程块中所创建的流启动的网格上执行同步。线程块只有等到其内所有线程启动的任务都完成后才能算完成，如果线程块中所有线程在所有的子任务完成之前要退出，同步操作会被自动触发

流和事件

CUDA流和事件允许控制网格启动间的依赖：被启动到同一条流的网格按序执行，事件可以被用来创建流之间的依赖。在设备上创建的流和事件都服务于这个相同的目的(控制子任务依赖)。在一个线程块中创建的流和事件存在于同一个线程块范围内，但是当在创建它们的线程块之外使用时会产生未知的行为。如上一小节所述，被一个线程块启动的任务在线程块退出时会发生隐式同步，包括启动到流中的任务，而且所有的依赖会被正确处置。超出线程块范围的流操作的行为结果是未知的，而且创建于主机的流和事件当在任何核函数中使用、创建于父网格中的流和事件在子网格中使用时也都会产生未知的行为。

顺序和并发

从设备运行时启动的核函数顺序服从CUDA流排序语义，在一个线程块内，所有启动到相同流中的核函数会按序执行；同一个块中的多个线程启动到同一流中时，流中的执行那个顺序取决于块内的线程调度机制，这可以通过同步原语控制，比如__syncthreads()。

注意，包括null流在内的流是被线程块中所有线程共享的。如果线程块中的多个线程启动到了隐式null流中，这些启动会按序执行，如果需要并发，就得用显式命名的流。

动态并行支持了程序内更容易被表达的并发，但是，设备运行时在CUDA执行模型里没有引入的新的并发担保，也就是不会担保设备上任何数量的不同线程块的并发执行。这种并发担保的缺乏也延伸到了父子网格中，当父线程块启动子网格时，子网格不会保证开始执行，直到父线程块到达了显式的同步点（比如cudaDeviceSynchronize()）。尽管并发经常被容易实现，但并发度也随着设备配置、应用负载和运行时调度这些函数的不同而不同，因此对不同线程块的任何并行度的依赖都是不安全的。

设备管理

在设备运行时没有多GPU的支持，运行时只有能力在它当前执行所在的设备上进行操作，但是允许查询系统上任何支持cuda的设备的属性。

内存模型

父子网格共享相同的全局和常量内存，但是有不同的局部和共享内存

连续性和一致性

全局内存：

父子网格拥有对全局内存的连续访问，并且父子之间会担保弱一致性。子网格执行中有两种情况可以和父网格的内存访问保持强一致性：子网格被其父调用或者子网格作为父线程调用的同步API的信号而完成时。在子网格调用之前进行的父线程所有内存操作都是对子网格可见的，父线程完成要在子网格完成点进行的同步后，子网格中的所有内存操作也是对其父可见的。

在下面的代码中，执行child_launch()的子网格只保证能看见在它启动前对data进行的修改。因为父线程的线程0执行子网格的启动，子网格就会和线程0可见的内存保持一致。由于第一个__syncthreads()调用，子网格会看到data[0] = 0, data[1] = 1, ....., data[255] = 255(否则子网格就只能看到data[0] = 0了)。开启cudaDeviceSynchronize()后，当子网格返回时，线程0才保证能看到子网格线程中做的修改，这些修改只会在第二个__syncthreads()调用后被父网格中的其他线程所见：

#include "cuda_runtime.h"
#include "stdio.h"

__device__ void printArray(int *data, int n) {
    for (int i = 0; i < n; i++) {
        printf("%d\t", data[i]);
    }
}

__global__ void child_launch(int *data, int n) {
    data[threadIdx.x] = data[threadIdx.x] + 1000;
}

__global__ void parent_launch(int *data, int n) {
    data[threadIdx.x] = threadIdx.x;

    __syncthreads();
    printArray(data, n);

    if (threadIdx.x == 0) {
        child_launch<<<1, 6>>>(data, n);
//        cudaDeviceSynchronize();
    }

    printArray(data, n);

    __syncthreads();
}

int main() {
    cudaSetDevice(0);

    int size = 6;
    int *h_data = (int *) malloc(size * sizeof(int));
    int *d_data;

    for (int i = 0; i < size; i++) {
        h_data[i] = 0;
    }
    cudaMalloc(&d_data, size * sizeof(int ));
    cudaMemcpy(d_data, h_data, size * sizeof(int ), cudaMemcpyHostToDevice);

    parent_launch<<<1, 6>>>(d_data, size);
    cudaDeviceSynchronize();

    cudaFree(d_data);
    free(h_data);
    return 0;
}

然后编译：

C:\Users\songzeceng\CLionProjects\CudaDemo>nvcc DynamicParallelTest.cu -gencode arch=compute_61,code=sm_61 -rdc=true -o DynamicParallelTest

运行即可，截图如下，可见由于线程0没有在子网格启动后调用同步函数，所以它的第二次输出还是6个0；而别的线程第二次输出都是在线程0调用同步函数__syncthreads()之后的，所以别的都能输出1000多的值

零复制内存：

零复制系统内存有着和全局内存一样的连续性和一致性保障，而且也服从上面讲的语义。核函数不允许分配或释放零复制内存，但是可以使用从主机传来的零复制内存指针。

常量内存：

常量不可修改，因此不能从设备、父子自动之间进行修改。也就是说，所有__constant__变量的值必须在启动核函数之前于主机中设置，常量内存自动从父核函数继承到子核函数中。从核函数线程中对常量内存变量取址有着和所有CUDA程序相同的语义，而且父子之间对此指针的传递是天然被支持的。

共享和局部内存：

共享和局部内存分别是线程块私有和线程私有的，因此在父子之间是不可见或者连续的。在定义域之外对共享或局部内存中的变量进行的操作会产生未定义的行为和错误。当英伟达编译器检测到指向局部或共享内存的指针作为参数传给核函数时，它不会尝试警告，在运行时，程序员可以使用__isGlobal()指令来确定一个指针是否引用了全局内存中的变量，从而可以安全地传给子启动。

注意调用cudaMemcpy*Async()或cudaMemset*Async()可能会在设备上调用新的子核函数以保留流语义，把局部或共享内存指针传给这种函数是非法的，并且会返回一个错误。

局部内存：

局部内存是执行线程私有的，在其外不可见。启动子核函数时把指向局部内存的指针当成参数传进去是非法的，而且从子核函数中析构这种内存指针会产生未定义的结果。比如下面的代码就是非法的，如果x_array被child_launch()访问还会产生未定义的结果：

int x_array[10];
child_launch<<<1, 1>>>(x_array);

有时程序员不容易意识到变量被编译器放到局部内存中了，不过一般的规则是：传给子核函数的所有变量应该被显式分配到全局内存堆的变量（通过cudaMalloc()、new()分配，或者是声明__device__的全局变量，如下所示）：

// Correct - "value" is global storage
__device__ int value;
__device__ void x() {
    value = 5;
    child<<< 1, 1 >>>(&value);
}

// Invalid - "value" is local storage
__device__ void y() {
    int value = 5;
    child<<< 1, 1 >>>(&value);
}

纹理内存：

相比纹理访问而言，向已经和纹理映射的全局内存区域写数据是不连续的。在调用子网格或子网格完成时，连续的纹理内存是硬性要求，这意味着在子核函数启动前的内存写会反映到子核函数的纹理内存访问中，类似地，子核函数的内存写也会被反应到父核函数的纹理内存访问中（在父核函数完成在子核函数完成后要进行的同步之后），父子核函数对内存的并发访问可能会产生不一致的数据

编程接口

CUDA C/C++引用

这一节描述为了支持动态并行而对CUDA C/C++语言拓展进行的改变和增补。为了支持动态并行，使用CUDA C/C++对CUDA核函数可用的语言接口(称之为设备运行时)和API基本和主机可用的CUDA运行时API类似。CUDA运行时API语法语义尽可能保留了下来，以便支持可以运行在主机或设备环境的程序进行简易的代码复用。与CUDA C/ c++中的所有代码一样，这里列出的api和代码都是线程代码，这些代码让每个线程可以做出特有的关于接下来执行什么核函数或者操作的动态决策。线程块内的线程执行任何设备运行时API没有同步要求，这样就可以让设备运行时API函数在任意分支的设备代码里被调用，而且没有死锁

设备端的核启动

核函数可以从设备端通过使用标准CUDA<<<>>>语法启动：kernel_name<<<Dg, Db, Ns, S>>>([kernel arguments]);，其中的参数描述见下表：

参数	类型	描述
Dg	dim3	网格的尺寸
Db	dim3	线程块的尺寸
Ns	size_t	本次调用除了静态分配的共享内存外，还要给每个线程块动态分配的共享内存字节数，默认为0
S	cudaStream_t	与本次调用相关的流，必须和当前调用使用的流一致，默认为0

启动是异步的：

和主机端的启动一样，所有的设备端核启动都是和启动线程异步的，也就是说<<<>>>启动命令会立刻返回，启动线程会继续执行，直到它碰到了名切的同步启动点，比如cudaDeviceSynchronize()。网格启动被发送到设备上，并且和父线程独立运行，子网格可以在启动后任何时间开始运行，但直到启动线程到达了明确的同步启动点，子网格才能保证开始执行。

启动环境配置：

所有全局的设备配置设置(比如cudaDeviceGetCacheConfig()返回的共享内存和L1缓存大小和cudaDeviceGetLimit()返回的设备限制等)都会从父单元那儿继承过来，也就是说如果当父单元启动时，全局的执行配置是16k的共享内存和48k的L1缓存，子单元的执行状态也会这么配置，同样诸如栈大小等设备限制的设置也会如法炮制。

对于主机启动的核函数，主机设置的每个核配置的优先级会覆盖全局设置，这些配置也会在从设备启动核函数时使用，从设备端从新配置核函数的环境是不可能的。

流

命名流和未命名(NULL)流都对设备运行时可用。命名流可以被线程块内任何线程使用，但是流句柄不能被传给其他线程块或者父子核函数。换句话说，流应该是创建它的块的私有物，流句柄不保证在线程块间是独一无二的，所以在没有分配此流句柄的线程块内使用流句柄，会导致未定义的行为。和主机端的启动类似，启动到不同的流上的工作可以并发运行，但实际的并发度无从保障。依赖于子核函数间并发度的程序不被CUDA编程模型支持，并且会导致未定义的行为。

主机端NULL流的跨流栅栏语义不在设备端支持，为了保持和主机端语义的兼容性，所有的设备流都必须使用cudaStreamCreateWithFlags()创建，传入cudaStreamNonBlocking标志，cudaStreamCreate()是只能在主机运行时使用的API，在设备上调用直接编译失败。cudaStreamSynchronize()和cudaStreamQuery()在设备运行时不可用，当应用需要知道流启动的子核函数已经完成时，cudaDeviceSynchronize()将会作为替代品来使用

隐式流(NULL流)：

在主机程序里，未命名(NULL)流有额外的栅栏同步语义。设备运行时提供了一个隐式的未命名流，在块内所有线程间共享，但因为所有的命名流必须使用cudaStreamNonBlocking标志创建，启动到NULL流中的任务不会和其他任何流中的挂起任务之间建立隐式依赖。

事件

只有CUDA事件的流间同步受到了支持，这意味着cudaStreamWaitEvent()被支持，但是cudaEventSynchronize()、cudaEventElapsedTime()和cudaEventQuery()都不被支持。因为cudaEventElapsedTime()不可用，就必须通过cudaEventCreateWithFlags()创建事件，并且要传入cudaEventDisableTiming标志。对于所有的设备运行时对象而言，事件对象可以在创建它们的线程块中的线程间共享，但是也拥有线程块本地性，不能被传给别的核函数或者一个核函数的别的块。事件句柄不能保证块间唯一，所以在没有分配此事件句柄的线程块内使用事件句柄，会导致未定义的行为。

同步

cudaDeviceSynchronize()函数会把线程块内任何线程启动的所有任务同步到此函数被调用的地方，注意这个函数也有可能在分支代码中被调用。如果调用线程想和从别的线程调用的子网格同步的话，程序就应该执行足够的额外线程间同步，比如通过调用__syncthreads()

块级别的同步：

cudaDeviceSynchronize()函数不意味着线程块内的同步，特别的，没有显式通过__syncthreads()函数进行的同步，调用线程不能对除了自己以外的其他线程启动了哪些工作做出假设。比如，如果一个块内的多个线程都在启动工作，并且需要为这些工作立马进行同步(可能因为基于事件的依赖)，那么应该由程序来保证这些工作在调用cudaDeviceSynchronize()函数之后由所有的线程提交。

因为这种实现允许对线程块内任何线程的任务启动进行同步，所以多个线程同时调用了cudaDeviceSynchronize()，那么可能只有第一次调用对所有启动的任务有效，剩下的调用就无效了，如下图所示

其中，只有t0的cudaDeviceSynchronize()起到了效果，而t1和t2的都被覆盖了

设备管理

只有运行核函数的设备被此核函数控制，这意味着诸如cudaSetDevice()这些API不在设备运行时上支持。从GPU上见到的活跃设备(由cudaGetDevice()返回)会有着和从主机系统中看到的相同的设备号。cudaDeviceGetAttribute()函数可以请求别的设备信息，因为这个函数允许指定设备id作为它的参数。注意，cudaGetDeviceProperties()这种一次获取所有的API函数不被设备运行时提供，属性必须挨个儿查询。

内存声明

设备和常量内存：

文件范围内用__device__或__constant__内存空间标识符声明的内存在和设备运行时有着同样的行为，所有核函数可以读写设备变量，不管此核函数的初始化启动是在主机运行时还是设备运行时完成的。同样的，对于模块范围内声明的__constant__s变量，所有的核函数都有相同的可见性。

纹理和表面：

CUDA支持动态创建纹理和表面对象(CUDA5.0开始)，此时纹理引用可以在主机上创建，传到核函数中并在里面使用，然后在主机上销毁。设备运行时不允许在设备嗲吗里创建销毁纹理或表面对象，但是从主机端创建的纹理和引用对象可以在设备里自由地使用和传。不管他们是在哪儿创建的，动态创建的纹理对象总是有效的，而且可以从父单元中传给子核函数。

注意，设备运行时不支持老的模块范围(比如Fermi风格的)的纹理和表面引用在设备启动的核函数里使用，它只能在主机中创建，但只能用在从主机启动的核函数里（也就是顶层核函数里）。

共享内存变量声明：

在CUDA C/C++里，共享内存既可以声明成静态大小的文件或函数范围变量也可以声明成由核函数调用器通过启动配置参数决定大小的外部(extern)变量，两种类型的声明都是设备运行时有效的，以下是后者的例子：

__global__ void permute(int n, int *data) {
    extern __shared__ int smem[]; // 由第三个启动参数决定大小
    if (n <= 1)
        return;

    smem[threadIdx.x] = data[threadIdx.x];
    __syncthreads();

    permute_data(smem, n);
    __syncthreads();

    // 不能把共享变量传给子核函数，所以要再写回全局内存
    data[threadIdx.x] = smem[threadIdx.x];
    __syncthreads();

    if (threadIdx.x == 0) {
        permute<<< 1, 256, n/2*sizeof(int) >>>(n/2, data);
        permute<<< 1, 256, n/2*sizeof(int) >>>(n/2, data+n/2);
    }
}

void host_launch(int *data) {
    permute<<< 1, 256, 256*sizeof(int) >>>(256, data);
}

符号地址：

就像核函数可见的地址空间内的所有全局设备变量一样，设备端符号(使用__device__标记的)可以在核函数内通过&操作符引用。这种方法也适用于__constant__符号，尽管这时指针指向的是只读数据。既然设备端的符号可以被直接引用，那些引用符号的CUDA运行时API(cudaMemcpyToSymbol()或cudaGetSymbolAddress()函数)就是多余的了，因此不被设备运行时支持。这意味着常量数据不能被正在运行的核函数改变，即便在启动子核函数之前，__constant__空间的引用也都是只读的。

API错误和启动失败

作为CUDA运行时的惯例，所有的函数都会返回错误代码。最后一个错误码被记录并且可以通过cudaGetLastError()函数获取，错误的记录是针对每个线程的，所以每个线程可以指明它遇到的最后的错误，错误码的类型是cudaError_t。和主机端启动类似，设备端的启动可能因为各种原因失败(比如无效参数等)，用户必须调用cudaGetLastError()来确定启动是否生成了一个错误，但是启动后没有错误不代表子核函数就能成功完成。

对于设备端异常，比如访问了无效的地址，子网格中的错误会被返回给主机，而不是父单元中的cudaDeviceSynchronize()调用。

启动设置API：

核函数是通过设备运行时库暴露的系统层面机制，因此可以通过cudaGetParameterBuffer()和cudaLaunchDevice()函数直接从ptx中获得。CUDA应用也可以自己调用这些函数，要求是和从ptx中调用一样的。对于这两种情况，用户要根据规范用正确的方式使用所有必要的数据结构，这些数据结构向后兼容。和主机端启动一样，设备端操作符<<<>>>会映射到指定的核函数启动API上，所以面向ptx的用户才可以执行核函数的启动，前端编译器也才能把<<<>>>翻译到这些调用上，下表是对cudaGetParameterBuffer()和cudaLaunchDevice()的说明：

运行时API启动函数	和主机运行时行为的区别(没有描述的地方就是没有区别)
cudaGetParameterBuffer	从<<<>>>中自动生成，注意和主机中对应函数的差别
cudaLaunchDevice	从<<<>>>中自动生成，注意和主机中对应函数的差别

这些启动函数的API和cuda运行时API中的有所不同，其中定义如下：

extern device cudaError_t cudaGetParameterBuffer(void **params);
extern __device__ cudaError_t cudaLaunchDevice(void *kernel, void *params, dim3 gridDim, dim3 blockDim, unsigned int sharedMemSize = 0, cudaStream_t stream = 0);

API参考

设备运行时支持的cuda运行时API在这里会详细列出。主机和设备运行时API有着同样的语法，除了列举出来的之外，语义也是一样的。下面的表格列举了这些API和主机端可用的对应版本的区别

ptx的设备端启动

本节是为面向并行线程执行(ptx)的编程语言和编译器、并且打算在他们的语言里支持动态并行的开发者而写的，提供了在ptx层面支持核启动的低层面细节

核启动API

设备端的核启动可以通过两个ptx可用的API来实现：cudaLaunchDevice()和cudaGetParameterBuffer()。前者使用通过调用后者得到参数缓存启动指定的核函数，并为这个核函数填充参数。如果不需要调用cudaGetParameterBuffer()，参数缓存就是null，并且启动的核函数没有任何参数

cudaLaunchDevice()：

cudaLaunchDevice()使用前，需要在在ptx层面用以下两种形式声明：

// 地址大小为64时的声明
.extern .func(.param .b32 func_retval0) cudaLaunchDevice
(
    .param .b64 func,
    .param .b64 parameterBuffer,
    .param .align 4 .b8 gridDimension[12],
    .param .align 4 .b8 blockDimension[12],
    .param .b32 sharedMemSize,
    .param .b64 stream
)
;

//  地址大小为32时的声明
.extern .func(.param .b32 func_retval0) cudaLaunchDevice
(
    .param .b32 func,
    .param .b32 parameterBuffer,
    .param .align 4 .b8 gridDimension[12],
    .param .align 4 .b8 blockDimension[12],
    .param .b32 sharedMemSize,
    .param .b32 stream
);

下面cuda层面的声明，会被映射到前述两种PTX层的声明之一：可以在系统头文件cuda_device_runtime_api.h。此函数在cudadevrt系统库中被定义，为了使用设备端核启动功能的话，应用必须和此库链接。

// CUDA-level declaration of cudaLaunchDevice()
extern "C" __device__
cudaError_t cudaLaunchDevice(void *func, void *parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize, cudaStream_t stream);

第一个参数是要被启动的核函数的指针，第二个参数是持有待启动核函数实际参数的参数缓存，其布局会在下一节讲到。其他参数就是启动的配置参数了，比如网格维度、线程块维度、共享内存大小和启动对应的流。

cudaGetParameterBuffer():

cudaGetParameterBuffer()也需要使用前在ptx层面中声明，声明方式也根据地址大小分为以下两种：

// 地址大小为64时的声明
.extern .func(.param .b64 func_retval0) cudaGetParameterBuffer
(
    .param .b64 alignment,
    .param .b64 size
);

// 地址大小为32时的声明
.extern .func(.param .b32 func_retval0) cudaGetParameterBuffer
(
    .param .b32 alignment,
    .param .b32 size
);

下面cuda层面的声明，会被映射到前述两种PTX层的声明之一：

// CUDA-level Declaration of cudaGetParameterBuffer()
extern "C" __device__
void *cudaGetParameterBuffer(size_t alignment, size_t size);

第一个参数指定了参数缓存的对齐要求，第二个参数是参数缓存字节数要求。在现有的实现中，被cudaGetParameterBuffer()函数返回的参数缓存总是可以保证64字节对齐，而第一个参数指定的对齐要求被忽略了。然而，还是建议传入正确的对齐要求参数(指定了被放到这个缓存中的任何参数的最大对齐情况)，来确保此函数在未来的可移植性。

参数缓存布局

参数缓存中的参数重排序是禁止的，而且每个放入到参数列表中的参数都要求是对齐的，也就是说，每个参数必须被放到参数缓存中的第n字节上，其中n是参数大小的最小整数倍，并且大于上一个参数所占的最后一个字节偏移量。比如上一个参数占到了第10个字节，而当前参数大小为6字节，那么当前参数就要被放到第12个字节上。参数缓存的最大大小为4KB。

关于被CUDA编译器生成的ptx代码的详细描述，请参见ptx-3.5规范

动态并行的工具包支持

CUDA代码里包含设备运行时API

和主机端的运行时API类似，CUDA设备运行时API的原型在程序编译阶段会被自动包含，因此不需要显式包含cuda_device_runtime_api.h

编译和链接

当使用nvcc编译时，CUDA程序自动和主机运行时库链接，但设备运行时是作为必须显式链接给要使用它的程序的E静态库被装载的。设备运行时被作为静态库提供(Windows中为cudadevrt.lib，Linux上为libcudadevrt.a)，想要使用设备运行时的GPU应用就必须链接它。链接设备库可以通过nvcc或nvlink完成，下面是简单的例子。

如果所有需要的源文件都可以从命令行指定的话，设备运行时程序可以被一步编译链接：

$ nvcc -arch=sm_61 -rdc=true hello_world.cu -o hello -lcudadevrt

也可以先把cuda的.cu源文件编译成目标文件，再把这些目标文件链接在一起：

$ nvcc -arch=sm_61 -dc hello_world.cu -o hello_world.o
$ nvcc -arch=sm_61 -rdc=true hello_world.o -o hello -lcudadevrt

可以参见The CUDA Driver Compiler NVCC的使用单独编译部分以查看更多细节

编程指南

基础

设备运行时是主机运行时的一个功能性子集，暴露了API层面的设备管理、核函数启动、设备内存复制、流管理和事件管理。编程已经体验过CUDA编程的人应该对设备运行时不陌生，设备运行时的语义语法和主机API非常相似，其差别也在本附录前文有所记载。下面的代码展示了一个使用动态并行的简单Hello World程序：

#include "cuda_runtime.h"
#include "stdio.h"

__global__ void child_launch() {
    printf("Hello ");
}

__global__ void parent_launch() {
    child_launch<<<1, 1>>>();
    cudaDeviceSynchronize();

    printf(" World!\n");
}

int main() {
    parent_launch<<<1, 1>>>();
    cudaDeviceSynchronize();

    return 0;
}

然后编译运行即可：

C:\Users\songzeceng\CLionProjects\CudaDemo>nvcc DynamicParallelTest.cu -gencode arch=compute_61,code=sm_61 -rdc=true -o DynamicParallelTest  && DynamicParallelTest.exe

输出结果不用说大家也知道：

性能

同步

一个线程的同步会影响到同一线程块内其他线程的表现，即便那些线程没有本身调用cudaDeviceSynchronize()，其影响取决于底层实现

支持动态并行的核函数负载

当控制动态启动时，活动的系统软件会影响同时运行的任何核函数的工作负载，不管系统软件自己是否调用了核函数。这种负载来自于设备运行时执行追踪和软件管理，相比主机端而言，这种情况可能会降低从设备端调用库的性能。通常来说，这种负载针对的是链接设备运行时库的应用。

实现限制和约束

动态并行会保证本附录描述的所有语义，然而，某些特定的硬件和软件资源是相互依赖实现的，由此限制了使用设备运行时应用的规模、性能和其他属性

内存足迹

设备运行时系统软件出于各种管理的目的会保留一些内存，主要是为了保存同步时父网格状态而做出的保留，其次时为了追踪挂起的网格启动而做出的保留。可以通过配置控制来减少这些预留区的大小以换取某些启动限制，具体请参见下文的配置选项。

预留内存的主要部分是作为父核函数状态的备份存储而分配的，当同步子网格启动时会使用。保守的说，这些内存必须支持存储设备上可能的最大数量的活跃内存的状态，这就意味着每次cudaDeviceSynchronize()调用时，根据设备配置，父辈都要牺牲150MB的内存，这对程序来说将是不可用的，即便这么大的内存没有完全地消耗。

嵌套和同步深度

使用设备运行时时，核函数可以启动另一个核函数，新的核函数可以再启动一个核函数，如此往复。每一次下级启动，我们称之为新的嵌套，嵌套层次的总数叫做程序的嵌套深度。同步深度被定义为子启动时程序需要显式同步的最深层次。具体来说，同步深度要比程序嵌套深度少一级(主机启动核函数时，程序嵌套深度为1，但此时没有子启动，所以没有同步深度)，但如果程序不需要在所有层调用cudaDeviceSynchronize()，那么同步深度可能和嵌套深度完全不同

最大的嵌套深度为24，但实际来说，真实的限制为系统对每一个新层的内存数量的要求（请参见上文内存足迹）。任何往比最大值更深一层的内核启动都会失败，注意这也适用于cudaMemcpyAsync()，这个函数可能自己会生成一个核启动，请参见配置选项一节。

默认情况下，充足的内存为两级同步所保留，最大同步深度(和预留存储)可以通过调用cudaDeviceSetLimit()函数并指定cudaLimitDevRuntimeSyncDepth来控制。被支持的层数必须在由主机启动的顶级核函数中配置，以保证嵌套程序的正确执行。在指定最大同步深度的更深一层调用cudaDeviceSynchronize()会返回一个错误。

当父核函数从没调用过cudaDeviceSynchronize()函数时，系统会发现不需要为父状态预留空间，这是一个优化点。在这种情况下，因为显式地父子同步从不发生，程序要求的内存足迹就会比保守最大值小得多，因此程序可以通过指定更浅的最大同步深度来避免后备内存的过度分配

挂起核函数启动

当核函数启动时，所有相关的配置和参数都会被追踪直到核函数完成，这些数据存在于一个系统管理的启动池中。这个启动池被划分为固定大小池和性能次之的虚拟池。设备运行时系统会先尝试在固定大小池中追踪启动数据，当此池满时，再使用虚拟化池追踪新的启动数据。固定大小的启动池的大小可以通过正在主机端调用cudaDeviceSetLimit()并指定cudaLimitDevRuntimePendingLaunchCount配置

配置选项

为设备运行时系统软件的资源分配可以通过在主机程序中调用cudaDeviceSetLimit()控制，这些限制必须在核函数启动前设置，当GPU正在运行程序时，就不能改变了。下面就是可以设置的限制：

限制	行为
cudaLimitDevRuntimeSyncDepth	设置cudaDeviceSynchronize()函数可以被调用的最大深度。核启动可以在比这个值更深的地方执行，但是在比这个值更深的地方进行显式同步会返回cudaErrorLaunchMaxDepthExceeded，此值默认为2
cudaLimitDevRuntimePendingLaunchCount	控制为缓存还没有开始执行的核启动而保留的内存数，会导致未解析的依赖或执行资源的缺乏。当缓存满了，设备运行时系统软件会尝试在性能稍次的虚拟缓存中追踪新的挂起启动。如果虚拟缓存也满了(比如所有可用的堆内存已经被消耗了)，启动不会发生，线程的最后一个错误会被设置成cudaErrorLaunchPendingCountExceeded，此值默认为2048个启动

内存分配和生命周期

cudaMalloc()和cudaFree()在主机和设备环境中有不同的语义。当从主机端调用时，cudaMalloc()会从会使用的设备内存中分配新的区域，当从设备运行时调用时，这些函数会映射到设备端的malloc()和free()，这意味着在设备环境下，总共可分配的内存受限于设备的malloc()函数所分配的堆大小(?在计算能力6.1下没得到验证)，这比可用的设备内存可能要小。而且，从主机系统调用cudaFree()来释放由设备的cudaMalloc()函数分配的指针会报错，反之亦然。

	主机端cudaMalloc()	设备端cudaMalloc()
主机端cudaFree()	支持	不支持
设备端cudaFree()	不支持	支持
分配限制	可用设备内存	cudaLimitMallocHeapSize