网站导航
> 文档中心 > CUDA异步并发之CUDA流详解

CUDA异步并发之CUDA流详解

网友: 文档中心 2022-05-30 22:14:35

CUDA中得异步并发

在这里插入图片描述

CUDA 将以下操作公开为可以彼此同时操作的独立任务：

  • 在主机上计算；
  • 设备上的计算；
  • 从主机到设备的内存传输；
  • 从设备到主机的内存传输；
  • 在给定设备的内存中进行内存传输；
  • 设备之间的内存传输。

这些操作之间实现的并发级别将取决于设备的功能和计算能力，如下所述。

主机和设备之间的并发执行

在设备完成请求的任务之前，异步库函数将控制权返回给宿主线程，从而促进了主机的并发执行。使用异步调用，许多设备操作可以在适当的设备资源可用时排队，由CUDA驱动程序执行。这减轻了主机线程管理设备的大部分责任，让它自由地执行其他任务。以下设备操作对主机是异步的:

  • 内核启动;
  • 内存复制在单个设备的内存中;
  • 从主机到设备内存拷贝的内存块大小不超过64kb的;
  • 由带有Async后缀的函数执行的内存拷贝;
  • 内存设置函数调用。
    程序员可以通过将CUDA_LAUNCH_BLOCKING环境变量设置为1来全局禁用系统上运行的所有CUDA应用程序的内核启动的异步性。此特性仅用于调试目的，不应用作使生产软件可靠运行的一种方法。

如果通过分析器（Nsight、Visual Profiler）收集硬件计数器，则内核启动是同步的，除非启用了并发内核分析。如果异步内存复制涉及非页面锁定的主机内存，它们也将是同步的。

并行执行内核

某些计算能力 2.x 及更高版本的设备可以同时执行多个内核。 应用程序可以通过检查 concurrentKernels 设备属性（请参阅设备枚举）来查询此功能，对于支持它的设备，该属性等于 1。

设备可以同时执行的内核启动的最大数量取决于其计算能力，并在表15 中列出。

来自一个 CUDA 上下文的内核不能与来自另一个 CUDA 上下文的内核同时执行。

使用许多纹理或大量本地内存的内核不太可能与其他内核同时执

数据传输和内核执行的重叠

一些设备可以在内核执行的同时执行与 GPU 之间的异步内存复制。 应用程序可以通过检查 asyncEngineCount 设备属性（请参阅设备枚举）来查询此功能，对于支持它的设备，该属性大于零。 如果复制中涉及主机内存，则它必须是页锁定的。

还可以与内核执行（在支持 concurrentKernels 设备属性的设备上）或与设备之间的拷贝（对于支持 asyncEngineCount 属性的设备）同时执行设备内复制。 使用标准内存复制功能启动设备内复制，目标地址和源地址位于同一设备上。

并行数据传输

某些计算能力为 2.x 及更高版本的设备可以重叠设备之间的数据拷贝。 应用程序可以通过检查 asyncEngineCount 设备属性（请参阅设备枚举）来查询此功能，对于支持它的设备，该属性等于 2。 为了重叠，传输中涉及的任何主机内存都必须是页面锁定的。

CUDA流

应用程序通过流管理上述并发操作。 流是按顺序执行的命令序列（可能由不同的主机线程发出）。 另一方面，不同的流可能会彼此乱序或同时执行它们的命令； 不能保证此行为，因此不应依赖其正确性（例如，内核间通信未定义）。 当满足命令的所有依赖项时，可以执行在流上发出的命令。 依赖关系可以是先前在同一流上启动的命令或来自其他流的依赖关系。 同步调用的成功完成保证了所有启动的命令都完成了。

创建与销毁

流是通过创建一个流对象并将其指定为一系列内核启动和主机 设备内存拷贝的流参数来定义的。 以下代码示例创建两个流并在锁页内存中分配一个浮点数组 hostPtr。

cudaStream_t stream[2];for (int i = 0; i < 2; ++i)    cudaStreamCreate(&stream[i]);float* hostPtr;cudaMallocHost(&hostPtr, 2 * size);

这些流中的每一个都由以下代码示例定义为从主机到设备的一次内存复制、一次内核启动和从设备到主机的一次内存复制的序列：

for (int i = 0; i < 2; ++i) {    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,      size, cudaMemcpyHostToDevice, stream[i]);    MyKernel <<>>   (outputDevPtr + i * size, inputDevPtr + i * size, size);    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,      size, cudaMemcpyDeviceToHost, stream[i]);}

每个流将其输入数组 hostPtr 的部分复制到设备内存中的数组 inputDevPtr，通过调用 MyKernel() 处理设备上的 inputDevPtr，并将结果 outputDevPtr 复制回 hostPtr 的同一部分。 重叠行为描述了此示例中的流如何根据设备的功能重叠。 请注意，hostPtr 必须指向锁页主机内存才能发生重叠。

通过调用 cudaStreamDestroy() 释放流:

for (int i = 0; i < 2; ++i)    cudaStreamDestroy(stream[i]);

如果调用 cudaStreamDestroy() 时设备仍在流中工作，则该函数将立即返回，并且一旦设备完成流中的所有工作，与流关联的资源将自动释放。

默认流

未指定任何流参数或等效地将流参数设置为零的内核启动和主机 设备内存拷贝将发布到默认流。因此它们按顺序执行。

对于使用 --default-stream per-thread 编译标志编译的代码（或在包含 CUDA 头文件（cuda.h 和 cuda_runtime.h）之前定义 CUDA_API_PER_THREAD_DEFAULT_STREAM 宏），默认流是常规流，并且每个主机线程有自己的默认流。

注意：当代码由 nvcc 编译时，#define CUDA_API_PER_THREAD_DEFAULT_STREAM 1 不能用于启用此行为，因为 nvcc 在翻译单元的顶部隐式包含 cuda_runtime.h。在这种情况下，需要使用 --default-stream 每个线程编译标志，或者需要使用 -DCUDA_API_PER_THREAD_DEFAULT_STREAM=1 编译器标志定义 CUDA_API_PER_THREAD_DEFAULT_STREAM 宏。

对于使用 --default-stream legacy 编译标志编译的代码，默认流是称为 NULL 流的特殊流，每个设备都有一个用于所有主机线程的 NULL 流。 NULL 流很特殊，因为它会导致隐式同步，如隐式同步中所述。

对于在没有指定 --default-stream 编译标志的情况下编译的代码， --default-stream legacy 被假定为默认值。

显式同步

有多种方法可以显式地同步流。

cudaDeviceSynchronize() 一直等待，直到所有主机线程的所有流中的所有先前命令都完成。

cudaStreamSynchronize() 将流作为参数并等待，直到给定流中的所有先前命令都已完成。 它可用于将主机与特定流同步，允许其他流继续在设备上执行。

cudaStreamWaitEvent() 将流和事件作为参数（有关事件的描述，请参阅事件），并在调用 cudaStreamWaitEvent() 后使添加到给定流的所有命令延迟执行，直到给定事件完成。

cudaStreamQuery() 为应用程序提供了一种方法来了解流中所有前面的命令是否已完成。

隐式同步

如果主机线程在它们之间发出以下任一操作，则来自不同流的两个命令不能同时运行：

  • 页面锁定的主机内存分配，
  • 设备内存分配，
  • 设备内存设置，
  • 两个地址之间的内存拷贝到同一设备内存，
  • 对 NULL 流的任何 CUDA 命令，
  • 计算能力 3.x 和计算能力 7.x 中描述的 L1/共享内存配置之间的切换。

对于支持并发内核执行且计算能力为 3.0 或更低的设备，任何需要依赖项检查以查看流内核启动是否完成的操作：

  • 仅当从 CUDA 上下文中的任何流启动的所有先前内核的所有线程块都已开始执行时，才能开始执行；
  • 阻止所有以后从 CUDA 上下文中的任何流启动内核，直到检查内核启动完成。

需要依赖检查的操作包括与正在检查的启动相同的流中的任何其他命令以及对该流的任何 cudaStreamQuery() 调用。 因此，应用程序应遵循以下准则来提高并发内核执行的潜力：

  • 所有独立操作都应该在依赖操作之前发出，
  • 任何类型的同步都应该尽可能地延迟。

重叠行为

两个流之间的执行重叠量取决于向每个流发出命令的顺序以及设备是否支持数据传输和内核执行的重叠（请参阅数据传输和内核执行的重叠）、并发内核执行（ 请参阅并发内核执行）和并发数据传输（请参阅并发数据传输）。

例如,在设备不支持并行数据传输,这两个流的代码示例创建和销毁不重叠,因为由stream[1]发起的内存复制会在stream[0]发起的内存复制之后执行。如果代码以以下方式重写(并且假设设备支持数据传输和内核执行的重叠)

for (int i = 0; i < 2; ++i)    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,      size, cudaMemcpyHostToDevice, stream[i]);for (int i = 0; i < 2; ++i)    MyKernel<<>>   (outputDevPtr + i * size, inputDevPtr + i * size, size);for (int i = 0; i < 2; ++i)    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,      size, cudaMemcpyDeviceToHost, stream[i]);

那么在stream[1]上从主机到设备的内存复制 与stream[0]上内核启动重叠。

在支持并发数据传输的设备上，Creation 和 Destruction 的代码示例的两个流确实重叠：在stream[1]上从主机到设备的内存复制 与在stream[0]上从设备到主机的内存复制甚至在stream[0]上内核启动（假设设备支持数据传输和内核执行的重叠）。但是，对于计算能力为 3.0 或更低的设备，内核执行不可能重叠，因为在stream[0]上从设备到主机的内存复制之后，第二次在stream[1]上内核启动，因此它被阻塞，直到根据隐式同步，在stream[0]上第一个内核启动已完成。如果代码如上重写，内核执行重叠（假设设备支持并发内核执行），因为在stream[0]上从设备到主机的内存复制之前，第二次在stream[1]上内核启动被。但是，在这种情况下，根据隐式同步，在stream[0]上从设备到主机的内存复制仅与在stream[1]上内核启动的最后一个线程块重叠，这只能代表总数的一小部分内核的执行时间。

Host函数(回调)

运行时提供了一种通过 cudaLaunchHostFunc() 在任何点将 CPU 函数调用插入到流中的方法。 在回调之前向流发出的所有命令都完成后，在主机上执行提供的函数。

以下代码示例在向每个流发出主机到设备内存副本、内核启动和设备到主机内存副本后，将主机函数 MyCallback 添加到两个流中的每一个。 每个设备到主机的内存复制完成后，该函数将在主机上开始执行。

void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void *data){    printf("Inside callback %d\n", (size_t)data);}...for (size_t i = 0; i < 2; ++i) {    cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice, stream[i]);    MyKernel<<>>(devPtrOut[i], devPtrIn[i], size);    cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost, stream[i]);    cudaLaunchHostFunc(stream[i], MyCallback, (void*)i);}

在主机函数之后在流中发出的命令不会在函数完成之前开始执行。

在流中的主机函数不得进行 CUDA API 调用（直接或间接），因为如果它进行这样的调用导致死锁，它可能最终会等待自身。

流优先级

可以在创建时使用 cudaStreamCreateWithPriority() 指定流的相对优先级。 可以使用 cudaDeviceGetStreamPriorityRange() 函数获得允许的优先级范围，按 [最高优先级，最低优先级] 排序。 在运行时，高优先级流中的待处理工作优先于低优先级流中的待处理工作。

以下代码示例获取当前设备允许的优先级范围，并创建具有最高和最低可用优先级的流。

// get the range of stream priorities for this deviceint priority_high, priority_low;cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);// create streams with highest and lowest available prioritiescudaStream_t st_high, st_low;cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high);cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low);
网络标签：

相关问题