CUDA设备运行时的编程分析

 CUDA设备运行时的编程分析

1. 基础知识

设备运行时是主机运行时的功能子集。API 级别的设备管理、内核启动、设备内存、流管理和事件管理从设备运行时公开。

已经有 CUDA 经验的人应该熟悉设备运行时的编程。设备运行时语法和语义与主机 API 的语法和语义大致相同，但本文档前面详述的任何例外情况除外。

以下示例显示了包含动态并行性的简单 Hello World 程序：

#include <stdio.h>

__global__ void childKernel()

{

    printf("Hello ");

}

__global__ void tailKernel()

{

    printf("World!\n");

}

__global__ void parentKernel()

{

    // launch child

    childKernel<<<1,1>>>();

    if (cudaSuccess != cudaGetLastError()) {

        return;

    }

    // launch tail into cudaStreamTailLaunch stream

    // implicitly synchronizes: waits for child to complete

    tailKernel<<<1,1,0,cudaStreamTailLaunch>>>();

}

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_75 -rdc=true hello_world.cu -o hello -lcudadevrt

2. 性能

2.1. 启用动态并行的内核开销

在控制动态启动时处于活动状态的系统软件可能会对当时正在运行的任何内核施加开销，无论它是否调用自己的内核启动。此开销来自设备运行时的执行跟踪和管理软件，并可能导致性能下降。通常，此开销是针对设备运行时库链接的应用程序产生的。

3. 实现限制和限制

动态并行性保证了本文档中描述的所有语义，但是，某些硬件和软件资源依赖于实现，并限制了使用设备运行时的程序的规模、性能和其他属性。

3.1. 运行时

3.1.1. 内存占用

设备运行时系统软件保留内存用于各种管理目的，特别是用于跟踪挂起的网格启动的保留。配置控件可用于减小此预留的大小，以换取某些启动限制。有关详细信息，请参阅下面的配置选项。

3.1.2. 挂起的内核启动

启动内核时，将跟踪所有相关的配置和参数数据，直到内核完成。此数据存储在系统管理的启动池中。

固定大小启动池的大小可通过从主机调用cudaDeviceSetLimit()并指定 cudaLimitDevRuntimePendingLaunchCount来配置。

3.1.3. 配置选项

设备运行时系统软件的资源分配通过cudaDeviceSetLimit()主机程序的 API 进行控制。必须在启动任何内核之前设置限制，并且在 GPU 主动运行程序时不得更改限制。

可以设置以下命名限制：

限制	行为
cudaLimitDevRuntimePendingLaunchCount	控制为缓冲由于未解析的依赖项或缺少执行资源而尚未开始执行的内核启动和事件预留的内存量。当缓冲区已满时，在设备端内核启动期间分配启动槽的尝试将失败并返回cudaErrorLaunchOutOfResources，而分配事件槽的尝试将失败并返回cudaErrorMemoryAllocation。启动槽的默认数量为 2048。应用程序可以通过设置 cudaLimitDevRuntimePendingLaunchCount来增加启动和/或事件槽的数量。分配的事件槽数是该限制值的两倍。
cudaLimitStackSize	控制每个 GPU 线程的堆栈大小（以字节为单位）。CUDA 驱动程序会根据需要自动增加每个内核启动的每线程堆栈大小。每次启动后，此大小不会重置回原始值。若要将每个线程的堆栈大小设置为其他值，可以调用 cudaDeviceSetLimit()来设置此限制。堆栈将立即调整大小，如有必要，设备将阻止，直到所有先前请求的任务完成。cudaDeviceGetLimit()可以调用以获取当前每线程堆栈大小。

3.1.4. 内存分配和生存期

cudaMalloc()与cudaFree()在主机和设备环境之间具有不同的语义。从主机调用时，从未使用的设备内存中cudaMalloc()分配一个新区域。从设备运行时调用时，这些函数映射到设备端malloc()和free() .这意味着在设备环境中，总可分配内存仅限于设备堆大小，该malloc()大小可能小于可用的未使用设备内存。此外，在设备上分配的指针上cudaMalloc()从主机程序调用cudaFree()是一个误差，反之亦然。

	cudaMalloc()在主机上	cudaMalloc()在设备上
cudaFree()在主机上	支持	不支持
cudaFree()在设备上	不支持	支持
分配限制	释放设备内存	cudaLimitMallocHeapSize

3.1.5. SM ID 和变形 ID

请注意，在 PTX 中%smid和%warpid定义为易失性值。设备运行时可能会将线程块重新调度到不同的 SM 上，以便更有效地管理资源。因此，在线程或线程块的整个生命周期内依赖%smid或%warpid保持不变是不安全的。

3.1.6. ECC 误差

没有 ECC 误差的通知可用于 CUDA 内核中的代码。整个启动树完成后，主机端将报告 ECC 误差。在执行嵌套程序期间出现的任何 ECC 误差都会生成异常或继续执行（取决于误差和配置）。