编程接口：
CUDA C / C ++参考：
本节介绍对用于支持动态并行性的CUDA C / C ++语言扩展的更改和补充。
使用CUDA C / C ++ for Dynamic Parallelism（称为设备运行时）的CUDA内核可用的语言界面和API基本上与主机上可用的CUDA运行时API类似。 尽可能保留CUDA Runtime API的语法和语义，以便于可以在主机或设备环境中运行的例程重用代码。
与CUDA C / C ++中的所有代码一样，这里列出的API和代码是每个线程代码。 这使得每个线程都能够针对接下来要执行的内核或操作做出独特的动态决策。 在一个块内的线程之间没有任何同步要求来执行任何提供的设备运行时API，这使得设备运行时API函数能够在没有死锁的情况下在任意发散的内核代码中被调用。
设备端内核启动：
内核可以使用标准的CUDA <<< >>>语法从设备启动：
kernel_name<<< Dg, Db, Ns, S >>>([kernel arguments]);- Dg是dim3类型，并指定网格的尺寸和大小
- Db类型为dim3，并指定每个线程块的大小和尺寸
- Ns的类型为size_t，并指定为此调用每个线程块动态分配的共享内存的字节数，并指定添加到静态分配的内存。 Ns是一个可选的参数，默认为0。
- S的类型为cudaStream_t，并指定与此调用关联的流。 该流必须已分配到正在进行调用的同一个线程块中。 S是一个可选参数，默认为0。
异步启动：
与主机端启动相同，所有设备端内核启动对于启动线程都是异步的。 也就是说，<<< >>>启动命令将立即返回，并且启动线程将继续执行，直到它达到显式启动同步点（如cudaDeviceSynchronize（））。 网格启动发布到设备，并将独立于父线程执行。 子网格可能在启动后的任何时间开始执行，但不保证在启动线程达到显式启动同步点之前开始执行。
启动环境配置：
所有全局设备配置设置（例如，从cudaDeviceGetCacheConfig（）返回的共享内存和L1高速缓存大小以及从cudaDeviceGetLimit（）返回的设备限制）将从父级继承。 也就是说，如果启动父代时，执行全局配置为共享内存16k和L1缓存48k，那么子执行状态将被配置为相同。 同样，设备限制（如堆栈大小）将保持配置。
对于主机启动的内核，从主机设置的每个内核配置将优先于全局设置。 内核从设备启动时也会使用这些配置。 无法从设备重新配置内核环境。
