CUDA C/C++ 优化数据传输的效率、流的使用

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

1、CUDA 流

CUDA中的流是按照主机代码发出的顺序在设备上执行的一系列操作。 虽然保证流中的操作按规定的顺序执行，但是可以交错不同流中的操作，并且在可能的情况下，它们甚至可以同时运行。默认流与其他流不同，因为它是关于设备上的操作的同步流：在设备上的任何流中的所有先前发出的操作完成之前，默认流中的任何操作都不会开始，并且在默认流中的操作必须在任何其他操作（在设备上的任何流中）开始之前完成。

2、默认流

CUDA中的所有设备操作（内核和数据传输）都在流中运行。 如果未指定任何流，则使用默认流（也称为“空流”）。请注意，2015年发布的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() 。

3、非默认流

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);

4、与流同步

由于非默认流中的所有操作对于主机代码都是非阻塞的，因此，可以将主机代码与流中的操作同步的情况下运行。

有几种方法可以做到这一点。其中一种方法是使用 cudaDeviceSynchronize() ，它会阻止主机代码，直到设备上以前发出的所有操作都完成。在大多数情况下，这种方法会因为整个设备和主机线程的暂停而真正影响性能。

CUDA流API具有多种不严格的方法来将主机与流同步。

函数 cudaStreamSynchronize(stream) 可用于阻止主机线程，直到指定流中的所有先前发出的操作完成为止。

函数 cudaStreamQuery(stream) 测试是否已完成向指定流发出的所有操作，而不会阻止主机执行。

函数 cudaEventSynchronize(event) 和 cudaEventQuery(event) 的行为与流的对应函数相似，不同之处在于它们的结果基于是否已记录指定的事件而不是指定的流是否空闲。

还可以使用 cudaStreamWaitEvent(event) 对特定事件在单个流中进行同步操作（即使该事件记录在其他流中或在其他设备上）。

5、内核执行与数据传输之间的重叠

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

（1）该设备必须能够“同时复制和执行”。 可以从 cudaDeviceProp 结构的 deviceOverlap 字段中查询，也可以从 CUDA SDK / Toolkit 附带的 deviceQuery 示例的输出中查询。 几乎所有具有计算功能1.1及更高版本的设备都具有此功能。

（2）内核执行和要重叠的数据传输都必须发生在不同的非默认流中。

（3）数据传输中涉及的主机内存必须是固定内存（pinned memory）。

因此，让我们修改上面的简单主机代码，使用多个流，看看是否可以实现任何重叠。在修改后的代码中，我们将大小为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,, 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, stream[i]);
}

上面显示的两种异步方法都能产生正确的结果，并且在两种情况下，依存操作都按照需要执行的顺序发给同一流。 但是，这两种方法的执行情况会有所不同，具体取决于所使用GPU的特定生成时间。

在GeForce GTX 1060上，计算能力为6.1，运行测试代码得出以下结：

Device : GeForce GTX 1060
Time for sequential transfer and execute (ms): 13.192288
max error: 1.192093e-07
Time for asynchronous V1 transfer and execute (ms): 5.993408
max error: 1.192093e-07
Time for asynchronous V2 transfer and execute (ms): 5.891360
max error: 1.192093e-07

完整的代码如下：

官方博客给的示例运行结果如下：

?这里第一次输出的是顺序传输和使用阻塞传输的内核执行，我们将其用作异步加速比较的基准。

为什么两种异步策略在不同的体系结构上表现不同？要了解这些结果，我们需要清楚地了解CUDA设备是如何安排和执行任务的。CUDA设备包含用于各种任务的引擎，这些引擎在发出操作时将操作排入队列。不同引擎中的任务之间的依赖关系得到维护，但在任何引擎中，所有外部依赖关系都会丢失；每个引擎队列中的任务都按其发出的顺序执行。

C1060具有单个复制引擎和单个内核引擎。 下图显示了在C1060上执行示例代码的时间线。

?在示意图中，我们假设主机到设备传输、内核执行和设备到主机传输所需的时间大致相同（选择内核代码是为了实现这一点）。正如顺序内核所期望的，任何操作都没有重叠。

对于我们代码的第一个异步版本，复制引擎中的执行顺序为：H2D流（1），D2H流（1），H2D流（2），D2H流（2），依此类推。

这就是为什么在C1060上使用第一个异步版本时，我们看不到任何加速的原因：任务以防止内核执行和数据传输重叠的顺序发布给复制引擎。

但是，对于版本2，如果所有主机到设备的传输都在任何设备到主机的传输之前发出，则重叠是可能的，执行时间较短。

从我们的示意图中，我们期望异步版本2的执行是顺序版本的8/12，即8.7 ms，这在前面给出的时序结果中得到了证实。

在C2050上，架构上的不同，因此与C1060有所不同。

C2050有两个拷贝引擎，一个用于主机到设备的传输，另一个用于设备到主机的传输，还有一个内核引擎。下图演示了我们的示例在C2050上的执行。

?具有两个复制引擎可以解释为什么异步版本1在C2050上可以实现良好的加速：流[i]中从设备到主机的数据传输不会阻止流[i + 1]中从主机到设备的数据传输 ，就像在C1060上所做的那样，因为C2050的每个复制方向都有一个单独的引擎。

该示意图预测执行时间相对于顺序版本将减少一半，这大致是我们的时序结果所显示的。

但是，在C2050的异步版本2中观察到的性能下降是为什么呢？ 这与C2050可以同时运行多个内核的能力有关。

当多个内核在不同（非默认）流中连续发出时，调度程序会尝试启用这些内核的并发执行，因此会延迟通常在每个内核完成（负责启动设备到主机的传输）之后出现的信号，直到所有内核完成。

因此，尽管在异步代码的第2版中主机到设备的传输和内核执行之间存在重叠，但是内核执行和设备到主机的传输之间没有重叠。

该示意图预测异步版本2的总时间为顺序版本的时间的9/12，即7.5 ms，这一点已由我们的计时结果证实。

好消息是，对于具有计算能力3.5（K20系列）的设备，Hyper-Q特性消除了定制启动顺序的需要，因此上述任何一种方法都将起作用。我们将在以后的文章中讨论使用 Kepler 特性，但现在，这里是在 Tesla K20c GPU 上运行示例代码的结果。这两种异步方法在同步代码上实现了相同的加速比。

?6、总结

本篇文章主要介绍流以及如何通过并发执行拷贝和内核来使用它们来覆盖数据传输时间。