在上一期的 C / C ++ 文章 中，我们讨论了如何在主机和设备之间高效地传输数据。在这篇文章中，我们讨论了如何将数据传输与主机上的计算、设备上的计算相重叠，在某些情况下，主机和设备之间的其他数据传输。实现数据传输和其他操作之间的重叠需要使用 CUDA 流，所以首先让我们了解一下流。
CUDA 流
CUDA 中的 stream 是按照主机代码发出的顺序在设备上执行的操作序列。虽然流中的操作被保证按规定的顺序执行，但是不同流中的操作可以被交错，并且在可能的情况下，它们甚至可以并发运行。
默认流
CUDA 中的所有设备操作（内核和数据传输）都在一个流中运行。如果没有指定流，则使用默认流（也称为“空流”）。默认流与其他流不同，因为它是关于设备上操作的同步流：在所有先前发出的操作 在设备上的任何流中 完成之前，默认流中的任何操作都不会开始，并且默认流中的操作必须在任何其他操作（在设备上的任何流中）之前完成就要开始了。
请注意， 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) 的行为与它们的流对应项相似，只是它们的结果基于是否记录了指定的事件，而不是基于指定的流是否空闲。您还可以使用 cudaStreamWaitEvent （ event ）在单个流中同步特定事件的操作（即使事件记录在不同的流中，或者记录在不同的设备上）。
重叠的内核执行和数据传输
前面我们演示了如何将默认流中的内核执行与主机上的代码执行重叠。但我们在这篇文章中的主要目标是向您展示如何将内核执行与数据传输重叠。要做到这一点有几个要求。
- 设备必须能够“并发复制和执行”。这可以从
cudaDeviceProp结构的deviceOverlap字段或从 CUDA SDK / Toolkit 附带的deviceQuery示例的输出中进行查询。几乎所有具有计算能力 1 . 1 及更高版本的设备都具有此功能。 - 要重叠的内核执行和数据传输必须同时发生在 different 、 non-default 流中。
- 数据传输所涉及的主机内存必须是 pinned 内存。
因此，让我们从上面修改我们的简单主机代码，以使用多个流，看看是否可以实现任何重叠。这个例子的完整代码是 在 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 ）上运行测试代码（来自 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 stream （ 1 ）、 D2H stream （ 1 ）、 H2D stream （ 2 ）、 D2H stream （ 2 ）等等。这就是为什么我们在 C1060 上使用第一个异步版本时看不到任何加速：任务是按照排除内核执行和数据传输重叠的顺序被发送到复制引擎的。然而，对于版本 2 ，在所有主机到设备的传输在任何设备到主机的传输之前发出，重叠是可能的，如较低的执行时间所示。根据我们的示意图，我们期望异步版本 2 的执行时间是顺序版本的 8 / 12 ，或者 8 . 7ms ，这在前面给出的计时结果中得到了确认。
在 C2050 上，两个功能相互作用导致与 C1060 不同的行为。 C2050 有两个复制引擎，一个用于主机到设备的传输，另一个用于设备到主机的传输，以及一个内核引擎。下图说明了我们的示例在 C2050 上的执行。
有两个复制引擎解释了为什么异步版本 1 在 C2050 上实现了很好的加速：流[i] 不阻止流中数据的主机到设备传输 [i + 1]中数据的主机到设备的传输，因为 C2050 上的每个复制方向都有一个单独的引擎。示意图预测了执行情况相对于顺序版本，时间被缩短一半，这大致就是我们的计时结果显示的。
但是在 C2050 上的异步版本 2 中观察到的性能下降呢？这与 C2050 并发运行多个内核的能力有关。当多个内核在不同（非默认）流中背靠背地发出时，调度程序尝试启用这些内核的并发执行，结果会延迟通常在每个内核完成后出现的信号（这负责启动设备到主机的传输），直到所有内核完成。因此，虽然在第二个版本的异步代码中，主机到设备的传输和内核的执行之间有重叠，但是内核执行和设备到主机的传输之间没有重叠。示意图预测异步版本 2 的总时间是顺序版本的 9 / 12 ，即 7 . 5 毫秒，这一点由我们的计时结果证实。
CUDA Fortran 异步数据传输 中提供了关于本文中使用的示例的更详细的描述，好消息是对于具有计算能力 3 . 5 （ K20 系列）的设备， Hyper-Q 特性消除了定制发布顺序的需要，因此上述任何一种方法都可以工作。我们将在以后的文章中讨论使用开普勒特性，但是现在，这里是在 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 专业提示： CUDA 7 流简化并发 ）。这在编写库时尤其重要。如果库中的代码使用默认流，那么最终用户就没有机会将数据传输与库内核执行重叠。
现在您已经知道如何在主机和设备之间高效地移动数据，所以我们将研究如何在 下一篇文章 中的内核中高效地访问数据。
