原文链接
https://developer.nvidia.com/blog/how-overlap-data-transfers-cuda-cc/
在我们上一篇 CUDA C/C++ 文章中，我们讨论了如何在主机和设备之间有效地传输数据。在这篇文章中我们将讨论数据传输与host端计算和device端计算如何进行交叠，并且在某些情况下，host端和device端之间的其他数据如何传输。实现数据传输和其他操作之间的重叠需要使用cuda流，所以首先让我们了解一下流
cuda流
CUDA 中的流是一系列操作，它们按照主机代码发出的顺序在设备上执行。虽然保证流中的操作按规定的顺序执行，但不同流中的操作可以交错，如果可能，它们甚至可以并发运行。
默认流
cuda中所有的device端操作（核函数与数据传输）都是运行在流中的。如果未指定流，则使用默认流（也称为“空流”）。默认流与其他流不同，因为它是与device上的操作相关的同步流：在device上任何默认流中先前发出的操作完成之前，默认流中的其他操作都不会开始，并且默认流中的操作必须在其他操作（device上的任何流）开始之前完成。
请注意，2015 年发布的 CUDA 7 引入了一个新选项，可以为每个host线程使用单独的默认流，并将每个线程的默认流视为常规流（即它们不与其他流中的操作同步）。在 GPU Pro Tip: CUDA 7 Streams Simplify Concurrency 一文中阅读有关此新行为的更多信息。（https://developer.nvidia.com/blog/gpu-pro-tip-cuda-7-streams-simplify-concurrency/）
让我们看一些使用默认流的简单代码示例，并从host和device的角度讨论操作如何进行。
cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice); increment<<<1,N>>>(d_a) cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);
在上面的代码中，从device的角度来看，所有三个操作都发布到同一个（默认）流，并将按照它们发布的顺序执行。
从host的角度来看，隐式数据传输是阻塞或同步传输，而内核启动是异步的。由于第一行host到device的数据传输是同步的，所以直到host到device的传输完成后，CPU 线程才会到达第二行的内核调用。内核开始执行后，CPU 线程移动到第三行，但由于device端的执行是顺序的（默认流），该行上的传输无法开始。从host角度来看，kernel启动的异步行为使得device计算和host计算之间的重叠变得非常简单。我们可以修改代码以添加一些独立的 CPU 计算，如下所示。
cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice); increment<<<1,N>>>(d_a) myCpuFunction(b) cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);
在上面的代码中，一旦 increment() 核函数在device上启动，CPU 线程就会执行 myCpuFunction()，将其在 CPU 上的执行与在 GPU 上的内核执行重叠。host函数还是device核函数先完成并不影响后续的device到host的数据传输，device端到host端的数据传输只有在核函数结束后才会开始。从device、的角度来看，与之前的示例没有任何变化；因为device端完全不知道 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)
流之间的同步
由于非默认流中的所有操作相对于host代码都是非阻塞的，因此会遇到需要将host代码与流中的操作同步的情况。有几种方法可以进行同步操作。比较笨重的方法是使用cudaDeviceSynchronize()，他会阻塞host‘代码，直到device上所有先前发出的操作完成。在大部分情况下，这种方法有点不值当，这会导致整个host和device的线程停顿而影响整个系统的性能。
cuda流的api有几种不太严格的同步host与流的方法。函数cudaStreamSynchronize(stream)可用于阻塞host线程，直到指定流中所有先前发出的操作完成。函数 cudaStreamQuery(stream) 测试发出到指定流的所有操作是否都已完成，而不会阻止主机执行。函数 cudaEventSynchronize(event) 和 cudaEventQuery(event) 的行为类似于它们的流对应项，不同之处在于它们的结果基于是否已记录指定的event而不是指定的流是否空闲。还可以使用 cudaStreamWaitEvent(event) 同步单个流中特定事件的操作（即使事件记录在不同的流中，或在不同的设备上！）。
核函数执行与数据传输的重叠
之前我们演示了如何将默认流中的内核执行与主机上的代码执行重叠。但我们在这篇文章中的主要目标是展示如何将内核执行与数据传输重叠。发生这种情况有几个要求。
- 设备必须能够“并发复制和执行”。这可以从 cudaDeviceProp 结构的 deviceOverlap 字段或从 CUDA SDK/工具包中包含的 deviceQuery 示例的输出中查询。几乎所有具有 1.1 及更高计算能力的设备都具有此功能。
- 内核执行和要重叠的数据传输都必须发生在不同的非默认流中。
- 参与数据传输的主机内存必须是pinned内存。
因此，让我们修改上面的简单host代码以使用多个流，看看我们是否可以实现任何重叠。完整代码在github（https://github.com/NVIDIA-developer-blog/code-samples/blob/master/series/cuda-cpp/overlap-data-transfers/async.cu）上可以获取。在修改后的代码中，我们将大小为 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]); }
另一种方法是将类似的操作批处理在一起，首先发出所有host到device的传输，然后是所有内核启动，然后是所有device到host的传输，如下面的代码所示。
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）上运行测试代码（来自 Github）给出以下结果。
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
这里第一次报告的是使用阻塞传输的顺序传输和内核执行，我们将其用作异步加速比较的基线。为什么这两种异步策略在不同的架构上表现不同？要解读这些结果，我们需要更多地了解 CUDA 设备如何调度和执行任务。CUDA 设备包含用于各种任务的引擎，这些引擎在操作发出时对其进行排队。不同引擎之间的任务存在依赖关系，但在一个单独的引擎内部，所有外部依赖关系都将丢失；每个引擎队列中的任务按照它们发出的顺序执行。C1060 有一个复制引擎和一个内核引擎。下图显示了在 C1060 上执行示例代码的时间线。
在示意图中，我们假设主机到设备传输、内核执行和设备到主机传输所需的时间大致相同（选择内核代码是为了实现这一点）。正如顺序内核所预期的那样，任何操作都没有重叠。对于我们代码的第一个异步版本，复制引擎中的执行顺序是：H2D 流（1）、D2H 流（1）、H2D 流（2）、D2H 流（2）等等。这就是我们在 C1060 上使用第一个异步版本时没有看到任何加速的原因：按顺序分配到拷贝引擎的任务在内核执行和数据传输上没有任何的重叠。然而，对于版本 2，所有host到device的传输任务都在device到host的传输任务发出之前发出，由于有重叠，执行时间较短。从我们的示意图中，我们预计异步版本 2 的执行时间是顺序版本的 8/12，即 8.7 毫秒，这在之前给出的计时结果中得到了确认。
在 C2050 上，两个功能相互作用导致与 C1060 的行为不同。C2050 有两个复制引擎，一个用于host到device传输，另一个用于device到host传输，以及一个内核引擎。下图说明了在 C2050 上执行我们的示例。
有两个复制引擎解释了为什么异步版本 1 在 C2050 上实现了良好的加速：流 i 中的device到host的数据传输不会像在 C1060 上那样阻止流 i 中的host到device的数据传输，因为C2050在每个复制方向都有一个单独的引擎。原理图预测执行时间相对于顺序版本将减少一半，这大致是我们的时序结果所显示的。
但是为什么在 C2050 上的异步版本 2 中观察到的性能下降呢？这与 C2050 同时运行多个内核的能力有关。当多个内核在不同（非默认）流中背靠背发出时，调度程序会尝试并发启动这些核函数，因此在所有内核完成前，每个核函数完成后发出的信号（负责device端到host端的传输）通常都会有所延迟。因此，虽然在我们的异步代码的第二个版本中，主机到设备传输和内核执行之间存在重叠，但内核执行和设备到主机传输之间没有重叠。该示意图预测异步版本 2 的总时间是顺序版本时间的 9/12，即 7.5 毫秒，我们的计时结果证实了这一点。
CUDA Fortran （http://www.pgroup.com/lit/articles/insider/v3n1a4.htm）异步数据传输中提供了对本文中使用的示例的更详细描述。好消息是，对于计算能力为 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
总结
这篇文章和上一篇文章讨论了如何优化主机和设备之间的数据传输。上一篇文章重点介绍了如何最大限度地减少执行此类传输的时间，而这篇文章介绍了流以及如何通过并发执行副本和内核来使用它们来屏蔽数据传输时间。
在处理流的帖子中，我应该提到使用默认流很方便开发代码——同步代码更简单——最终您的代码应该使用非默认流或 CUDA 7 支持每线程默认流（阅读 GPU Pro Tip: CUDA 7 Streams Simplify Concurrency）。这在编写库时尤其重要。如果库中的代码使用默认流，则最终用户没有机会将数据传输与库内核执行重叠。
现在你知道如何在host和device之间有效地移动数据，所以我们将在下一篇文章中研究如何从内核中有效地访问数据。