工具包支持动态并行：
在CUDA代码中包含设备运行时API：
与主机端运行时API类似，CUDA设备运行时API的原型在程序编译期间自动包含在内。 没有必要明确包含cuda_device_runtime_api.h。
编译和链接：
在使用nvcc编译时，CUDA程序会自动与主机运行时库链接，但设备运行时作为静态库提供，必须明确地将其与希望使用它的程序链接。
设备运行时以静态库的形式提供（Windows上的cudadevrt.lib，Linux和MacOS下的libcudadevrt.a），使用设备运行时的GPU应用程序必须与之链接。 设备库的链接可以通过nvcc和/或nvlink完成。 下面显示了两个简单的例子。
如果可以从命令行指定所有必需的源文件，则可以在一个步骤中编译和链接设备运行时程序：

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

也可以先将CUDA .cu源文件编译为目标文件，然后将这些文件连接在一起分为两个阶段：

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

请参阅CUDA驱动程序编译器NVCC指南的使用单独编译部分了解更多详细信息。
编程指南：
设备运行时间是主机运行时间的一个功能子集。 API级设备管理，内核启动，设备memcpy，流管理和事件管理从设备运行时暴露出来。
对于已经具有CUDA经验的人员，设备运行时编程应该很熟悉。 设备运行时语法和语义与主机API的语法和语义基本相同，本文档前面详细介绍了任何例外情况。
以下示例显示了一个包含动态并行性的简单Hello World程序：

#include <stdio.h>
__global__ void childKernel()
{
    printf("Hello ");
}
__global__ void parentKernel()
{
    // launch child
    childKernel << <1, 1 >> >();
    if (cudaSuccess != cudaGetLastError()) {
        return;
    }
    // wait for child to complete
    if (cudaSuccess != cudaDeviceSynchronize()) {
        return;
    }
    printf("World!\n");
}
int main(int argc, char *argv[])
{
    // launch parent
    parentKernel << <1, 1 >> >();
    if (cudaSuccess != cudaGetLastError()) {
        return 1;
    }
    // wait for parent to complete
    if (cudaSuccess != cudaDeviceSynchronize()) {
        return 2;
    }
    return 0;
}

这个程序可以在命令行的单个步骤中构建，如下所示：

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

性能:
同步:
即使其他线程本身不调用cudaDeviceSynchronize（），一个线程同步也会影响同一个线程块中其他线程的性能。 这种影响将取决于基础实施。
支持动态并行机制的内核开销:
在控制动态启动时处于活动状态的系统软件可能会对当时正在运行的任何内核施加开销，无论它是否调用其自己的内核启动。 这种开销是由设备运行时的执行跟踪和管理软件产生的，并且与从主机侧相比，可能导致例如由设备进行库调用时的性能下降。 总的来说，这种开销是针对与设备运行时库链接的应用程序而产生的。