本文翻译自NVIDIA官方博客Parallel Forall，内容仅供参考，如有疑问请访问原网站:https://devblogs.nvidia.com/parallelforall/how-overlap-data-transfers-cuda-cc/

在上一篇博客中，我们讨论了如何在主机和设备之间高效地进行数据传输。在这篇文章中，我们将讨论如何使用主机端的计算、设备端的计算以及某些情况下的主机与设备端的数据传输来隐藏数据传输。要实现使用其他操作隐藏数据传输需要使用CUDA流，所以首先让我们来了解一下CUDA流。

译者注：这里为了符合中文的习惯，我将“Overlap Data Transfers”译为“隐藏数据传输”。“overlap”，原意为重叠，这里将其翻译为隐藏，既可以表达隐藏了数据传输的开销，也可以隐含地表达重叠的意思，更加的形象贴切。但是某些地方，为了表达顺畅，我也将其直接翻译为重叠。不管翻译成什么，只需要明白隐藏就是靠重叠来实现的，通过将几种相同或不同的操作重叠，我们就可以近似地实现隐藏某些开销。

CUDA流

CUDA流是由主机端发布，在设备端顺序执行的一系列操作。在一个CUDA流中的操作可以保证按既定的顺序执行，而在不同的流中的操作可以交叠执行，有时甚至可以并发(concurrently)执行。

默认流

所有设备操作，包括核函数和数据传输，都运行在CUDA流中。当没有指定使用哪个流时，就会使用默认流(也叫做“空流”，null stream)。默认流不同于其他流，因为它是一个对于设备上操作同步的CUDA流：直到之前发布在流中的所有操作完成，默认流中的操作才会开始；默认流中的操作必须在其他流中的操作开始前完成。

请注意在2015年发布的CUDA 7引入了一个新的特性——可以在每个主机线程中使用单独的默认流；也可以将每个线程的默认流作为普通流使用(即它们不对其他流中的操作进行同步)。详情请阅读这篇文章——GPU Pro Tip: CUDA 7 Streams Simplify Concurrency。

让我们来一起看一个使用默认流的简单例子，以及讨论如何从主机和设备的角度分析流中操作的执行过程。

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()，这样就实现了主机端myCpuFunction执行与设备端核函数执行的重叠。无论是主机端的函数先执行还是设备端的核函数先执行都不会影响之后设备到主机的数据传输，因为只有在核函数执行完毕之后它才会开始。从设备的角度来看，与前一个代码相比什么也没有改变，设备完全不会意识到myCpuFunction()的执行。

非默认流

非默认流在主机端声明、创建、销毁的C/C++代码如下：

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

我们可以使用cudaMemcpyAsync()函数来在一个非默认流中发布一个数据传输，这很类似于之前博客中讨论的cudaMemcpy()函数，区别就在于前者有第四个参数，用于标识使用哪个CUDA流。

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

cudaMemcpyAsync()在主机端是非同步的，所以当数据传输一旦开始控制权就会立刻返回到主机线程。对于2D和3D的数组的拷贝，我么可以使用cudaMemcpy2DAsync()和cudaMemcpy3DAsync()的函数形式。

在启动核函数时，我们需要使用第四个执行时配置参数(三对尖括号中)——流标识符(第三个执行时配置参数是为了分配共享内存，我们会在之后讨论，这里使用0)。

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

流的同步

你可能会遇到需要将主机代码与流中操作同步的情况，但是非默认流中的所有操作对于主机代码都是非同步的。有好几种方法可以解决这个问题。最有力的方法是使用cudaDeviceSynchronize()，它可以阻塞主机代码直到之前所有发布在设备端的代码全部完成为止。在大多数情况下，这其实都太过了，而且也会有损程序性能，因为这种方式会拖延整个设备和主机线程。

译者注：流的同步一般被用于时间测量。

CUDA流API中有多种温和的方式来同步主机代码。函数cudaStreamSynchronize(流)可以用于阻塞主机线程直到之前发布在指定流的所有操作完成为止。函数cudaStreamQuery(流)可以用于测试之前发布在指定流的所有操作是否完成，但不会阻塞主机线程。函数cudaEventSynchronize(事件)和cudaEventQuery(事件)与前两种函数很像，区别在于后者是基于指定事件是否被记录而前者是基于指定的流是否空闲。你也可以在一个单独的流中基于一个特定的事件使用cudaStreamWaitEvent(事件)函数(即使事件被记录在不同的流中或者不同的设备中！)

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

之前我们已经演示了如何在默认流中用主机端代码来隐藏核函数执行。但是我们的主要目的是演示如何用核函数执行隐藏数据传输。要实现它有几点要求：

• 设备必须可以“并发地拷贝和执行”。我们可以通过访问cudaDeviceProp结构体的deviceOverlap属性或者从CUDA SDK/Toolkit中deviceQuery示例程序的输出中获得。几乎所有计算能力1.1及以上的设备都支持设备重叠。

• 核函数执行和数据传输必须在不同的非默认流中。

• 涉及到数据传输的主机内存必须是固定主机内存。

下面让我们来修改上面的代码以使用多个CUDA流，看一看是否实现了数据传输的隐藏。完整的代码可以在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的GPU(计算能力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

这里数据传输和核函数顺序执行的同步版本可以作为比较上述两种异步版本是否有加速效果的基准。为什么这两种异步执行策略在不同架构上的效果不同呢？为了解释这一结果，我们需要了解CUDA设备如何调度和执行任务。CUDA设备中存在多种不同任务的引擎，它们会对发布的操作进行排队。它们的功能就是维护不同引擎中任务间的依赖，但是在引擎内部所有的外部依赖都会丢失；每个引擎中的任务都会按照它们被发布的顺序执行。C1060有一个单独的拷贝引擎和一个单独的核函数引擎。下图是C1060运行上面示例代码的时间线：

C1060 timelines

NOTE:H2D表示主机到设备；D2H表示设备到主机

在这个原理图中，我们假设主机到设备的数据传输、核函数执行、设备到主机三者所用的时间相同(所选择的核函数代码就是专门这样设计的)。正如预料的那样，顺序执行的核函数并没有任何操作重叠。对于异步版本1的代码，拷贝引擎中的执行顺序是： H2D 1号流, D2H 1号流, H2D 2号流, D2H 2号流, 以此类推。这就是为什么异步版本1没有任何加速的原因：在拷贝引擎上任务的发布顺序使得核函数执行和数据传输无法重叠。然而，从版本2较少的执行时间来看，所有主机到设备的数据传输都在设备到主机的数据传输之前，是有可能实现重叠的。在原理图中，我们可以看出异步版本理论时间是顺序版本的8/12，前面的结果8.7ms刚好符合这个推算。

在C2050中，有两个特征共同导致了它与C1060的性能差异。C2050有两个拷贝引擎，一个是用于主机到设备的数据传输，另一个用于设备到主机的数据传输，第三个引擎是核函数引擎。下图描述了C2050执行示例代码的时间线：

c2050 timelines

C2050具有两个拷贝引擎恰好解释了为什么异步版本1在C2050上具有很好的加速效果：与C1060正相反，在stream[i]上设备到主机的数据传输并不会妨碍stream[i+1]上的主机到设备的数据传输，因为在C2050上每个方向的拷贝都有单独的引擎。上面的原理图显示，该异步版本1的执行时间大约是顺序版本的一半，和实际结果相差无几。

但是我们该如何解释异步版本2在C2050上的性能下降呢？其实这与C2050可以并发执行多个核函数有关。当多个核函数背靠背地被发布在不同的流(非默认流)中时，调度器会尽力确保这些核函数并发执行，结果就导致每个核函数完成的信号被延迟，即所有核函数执行完毕才发出信号，而这个信号负责启动设备到主机的数据传输。因此，在异步版本2中，主机到设备的数据传输与核函数执行可以重叠，而核函数执行与设备到主机的数据传输不能重叠。上面的原理图中显示异步版本2的总体时间大约是顺序版本的9/12，正好与实验结果7.5ms相吻合。

关于这个例子，在这篇文章CUDA Fortran Asynchronous Data Transfers中有更详细的讲解。让人高兴的是，对于计算能力3.5的设备(K20系列)，它所具有的超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支持的每个线程单独的默认流。尤其是在写库函数时，这一点尤为重要。如果在库函数中使用默认流，那么对于库函数用户就不会有机会实现数据传输和核函数执行的重叠了。

现在你应该明白了如何高效地在主机和设备间传输数据，在下一篇博客中我们开始学习如何在核函数中高效的访问数据。