CUDA - 如何在CUDA C/C++中重叠数据传输

pla66

已于 2023-06-11 23:59:11 修改

阅读量349

点赞数 1

分类专栏： CUDA 文章标签： c++ c语言

于 2023-06-02 16:45:58 首次发布

原文链接：https://developer.nvidia.com/blog/how-overlap-data-transfers-cuda-cc/

版权

CUDA 专栏收录该内容

8 篇文章 2 订阅

订阅专栏

原文链接：How to Overlap Data Transfers in CUDA C/C++

文章目录

在上一篇CUDA C/C++文章中，我们讨论了如何在主机和设备之间高效地传输数据。在本文中，我们讨论如何将数据传输与主机上的计算、设备上的计算以及在某些情况下主机和设备之间的其他数据传输重叠。实现数据传输和其他操作之间的重叠需要使用CUDA流（stream），因此首先让我们了解流。

CUDA流

CUDA中的流是一系列按照主机代码发出的顺序在设备上执行的操作。一个流内的操作可以保证按顺序执行，但不同流中的操作可以交错执行甚至并发执行。

默认流

CUDA中的所有设备操作（内核和数据传输）都在流中运行。如果没有指定流，则使用默认流（也称为“空流(null stream)”）。默认流与其他流不同，因为它是一个与设备上的操作相关的同步流：在设备上的任何流中所有先前发布的操作完成之前，默认流中的任何操作都不会开始；在开始（在设备上任何流中的）任何其他操作之前，必须完成默认流中的操作。

请注意，2015年发布的CUDA 7引入了一个新选项，即每个主机线程使用一个单独的默认流，并将每个线程的默认流视为常规流（即它们不与其他流中的操作同步）。了解更多：GPU专业提示：CUDA 7流简化并发。

让我们看一些使用默认流的简单代码示例，并从主机和设备的角度讨论操作是如何进行的。

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a)
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

在上面的代码中，从设备的角度来看，所有三个操作都被发布到同一个（默认）流，并将按照发布的顺序执行。

从主机的角度来看，隐式数据传输是阻塞或同步传输，而内核启动是异步的。由于第一行的主机到设备数据传输是同步的，因此CPU线程在主机到设备传输完成之前不会到达第二行上的内核调用。当内核被启动，CPU线程就会移动到第三行，但由于设备端的执行顺序，该行上的传输无法开始。

从主机的角度来看，内核启动的异步行为使得将设备和主机计算重叠变得非常简单。我们可以修改代码以添加一些独立的CPU计算：

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a)
myCpuFunction(b)
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

在上面的代码中，只要increment()核函数在设备上启动，CPU线程就会执行myCpuFunction()，从而使CPU上的执行与GPU上的内核执行重叠。无论是主机函数还是设备内核首先完成，都不会影响后续的设备到主机的传输，此传输只有在内核完成后才会开始。从设备的角度来看，与之前的示例相比没有任何变化；设备完全不知道myCpuFunction()。

非默认流

CUDA C/C++中的非默认流在主机代码中按以下方式声明、创建和销毁：

cudaStream_t stream1;
cudaError_t result;
result = cudaStreamCreate(&stream1);
result = cudaStreamDestroy(stream1);

为了向非默认流发出数据传输，我们使用了cudaMemcpyAsync()函数，该函数与前面的文章中讨论的cudaMemcpy()函数类似，但将流标识符作为第五个参数。

result = cudaMemcpyAsync(d_a, a, N, cudaMemcpyHostToDevice, stream1);

cudaMemcpyAsync()在主机上是非阻塞的，因此在发出传输操作后，控制权立即返回到主机线程。此例程有cudaMemcpy2DAsync()和cudaMemcpy3DAsync()变体，它们可以在指定的流中异步传输2D和3D数组。

为了将内核发布到非默认流，我们将流标识符作为第四个内核执行配置参数（第三个执行配置参数用于分配共享设备内存，我们稍后将讨论，现在使用0即可）。

increment<<<1,N,0,stream1>>>(d_a);

流同步

由于非默认流中的所有操作相对于主机代码都是非阻塞的，因此您将遇到需要将主机代码与流中的操作同步的情况。有以下几种方法：“粗暴”的方法是使用cudaDeviceSynchronize()，它会阻塞主机代码，直到之前在设备上发布的所有操作都完成为止。在大多数情况下这么做用力过猛，并且可能会因为整个设备和主机线程的停滞而损害性能。

CUDA流API有多种较轻的方法来同步主机和流。函数cudaStreamSynchronize(stream)可用于阻塞主机线程，直到指定流中所有先前发出的操作完成。函数cudaStreamQuery(stream)测试指定流中发出的所有操作是否已完成，同时不会阻止主机执行。函数cudaEventSynchronize(event)和cudaEventQuery(event)的作用与对应的stream函数类似，只是它们的结果是基于是否记录了指定的事件，而不是指定的流是否空闲。您还可以在一个流内使用cudaStreamWaitEvent(event)针对一个特定事件进行同步操作（即使该事件记录在不同的流中或不同的设备上！）。

重叠内核执行和数据传输

前面我们演示了如何将默认流中的内核执行与主机上的代码执行重叠。但这篇文章中的主要目标是向您展示如何将内核执行与数据传输重叠。要做到这一点，有几个要求：

该设备必须能够“并行复制和执行（concurrent copy and execution）”。这可以通过cudaDeviceProp结构体的deviceOverlap字段查询，也可以从CUDA SDK/Toolkit中包含的deviceQuery示例的输出中查询。几乎所有具有1.1及更高计算能力的设备都具有这种能力。
内核执行和要重叠的数据传输必须同时发生在不同的非默认流中。
数据传输中涉及的主机内存必须是固定内存（pinned memory）。

下面让我们修改上文中的简单主机代码，使用多个流，看看我们能否实现任何重叠。这个例子的完整代码可以在Github上找到。在修改后的代码中，我们将大小为N的数组分解为大小为streamSize的部分构成的块。由于内核在所有部分上独立运行，因此每个块都可以独立处理。使用的（非默认）流的数量为nStreams=N/streamSize。有多种方法可以实现数据的域分解和处理；一种是循环遍历数组中每个块的所有操作，如本示例代码所示。

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  cudaMemcpyAsync(&d_a[offset], &a[offset], streamBytes, cudaMemcpyHostToDevice, stream[i]);
  kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
  cudaMemcpyAsync(&a[offset], &d_a[offset], streamBytes, cudaMemcpyDeviceToHost, stream[i]);
}

另一种方法是将类似的操作一起批处理，首先发出所有主机到设备的传输，然后是所有内核启动，然后是所有设备到主机的传输，如以下代码所示。

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  cudaMemcpyAsync(&d_a[offset], &a[offset], 
                  streamBytes, cudaMemcpyHostToDevice, cudaMemcpyHostToDevice, stream[i]);
}

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
}

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  cudaMemcpyAsync(&a[offset], &d_a[offset], 
                  streamBytes, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToHost, stream[i]);
}

上面显示的两种异步方法都会产生正确的结果，而且在这两种情况下，依赖操作都会按照需要执行的顺序发布到同一个流。但根据使用的GPU版本不同，这两种方法的表现非常不同。在Tesla C1060（计算能力1.3）上运行测试代码，得到以下结果。

Device : Tesla C1060

Time for sequential transfer and execute (ms ): 12.92381
  max error : 2.3841858E -07
Time for asynchronous V1 transfer and execute (ms ): 13.63690 
  max error : 2.3841858E -07
Time for asynchronous V2 transfer and execute (ms ): 8.84588
  max error : 2.3841858E -07

在Tesla C2050（计算能力2.0）上得到以下结果。

Device : Tesla C2050

Time for sequential transfer and execute (ms ): 9.984512
  max error : 1.1920929e -07
Time for asynchronous V1 transfer and execute (ms ): 5.735584 
  max error : 1.1920929e -07
Time for asynchronous V2 transfer and execute (ms ): 7.597984
  max error : 1.1920929e -07

这里第一次报告的是使用阻塞传输的顺序传输和内核执行的用时，我们将其用作异步加速的比较基准。为什么这两种异步策略在不同的GPU架构上表现不同？为了解读这些结果，我们需要更多地了解CUDA设备如何调度和执行任务。CUDA设备包含用于处理各种任务的引擎，这些引擎在操作发出时对操作进行排队。不同引擎中的任务之间的依赖关系得到了维护，但在任何引擎内部，所有外部依赖关系都会丢失；每个引擎队列中的任务都是按照发布顺序执行的。C1060有一个复制引擎和一个内核引擎。下图显示了在C1060上执行示例代码的时间线。
在这里插入图片描述
在示意图中，我们假设主机到设备传输、内核执行和设备到主机传输所需的时间大致相同（内核代码经过特殊设计以实现这一点）。正如顺序内核所预期的那样，在任何操作中都没有重叠。对于代码的第一个异步版本，在复制引擎中执行的顺序是：H2D stream(1), D2H stream(1), H2D stream(2), D2H stream(2)，以此类推。这就是为什么我们在C1060上使用第一个异步版本代码时没有看到任何加速：任务以排除了内核执行和数据传输之间的任何重叠的顺序发布给复制引擎。然而，对于版本二，其中所有主机到设备的传输都是在设备到主机的传输之前发出的，这样重叠是可能的，如较低的执行时间所示。根据我们的示意图，我们预计异步代码版本2的执行是顺序版本的8/12，或8.7ms，这在前面给出的定时结果中得到了证实。

在C2050上，有两个特性相互作用，导致与C1060的行为不同。C2050有两个复制引擎，一个用于主机到设备的传输，另一个用于设备到主机的传输，还有一个内核引擎。下图说明了我们在C2050上执行的示例。
在这里插入图片描述
拥有两个复制引擎可以解释为什么异步代码版本1在C2050上实现了良好的加速：stream[i]中的设备到主机的数据传输不会像在C1060上那样阻塞stream[i+1]中的主机到设备的数据传输，因为C2050上的每个复制方向都有一个单独的引擎。该示意图预测，相对于顺序版本，执行时间将减半，这与计时结果大致相同。

但是在C2050上的异步版本2中观察到的性能下降怎么解释呢？这与C2050同时运行多个内核的能力有关。当多个内核在不同的（非默认）流中背靠背发布时，调度器试图启用这些内核的并发执行，因此会将通常在每个内核完成后出现的信号（负责启动设备到主机的传输）延迟处理，直到所有内核完成。因此，尽管在我们的异步代码的第二个版本中，主机到设备的传输和内核执行之间存在重叠，但内核执行和设备到主机的传输之间没有重叠。该示意图预测异步版本2的总时间为顺序版本时间的9/12，即7.5ms，计时结果证实了这一点。

CUDA Fortran异步数据传输中提供了对本文中使用的示例的更详细描述。好消息是，对于具有计算能力3.5的设备（K20系列）来说，其Hyper-Q功能无需特定的发布顺序，因此上述任何一种方法都适用。我们将在未来的文章中讨论使用Kepler功能，但目前，以下是在Tesla K20c GPU上运行示例代码的结果。正如您所看到的，两种异步方法相比于同步代码实现了相同的加速效果。

Device : Tesla K20c
Time for sequential transfer and execute (ms): 7.101760
  max error : 1.1920929e -07
Time for asynchronous V1 transfer and execute (ms): 3.974144 
  max error : 1.1920929e -07
Time for asynchronous V2 transfer and execute (ms): 3.967616 
  max error : 1.1920929e -07