[译]在CUDA C/C++中如何隐藏数据传输

本文翻译自NVIDIA官方博客Parallel Forall，内容仅供参考，若有疑问请访问原网站:https://devblogs.nvidia.com/p...

在上一篇博客中，咱们讨论了如何在主机和设备之间高效地进行数据传输。在这篇文章中，咱们将讨论如何使用主机端的计算、设备端的计算以及某些状况下的主机与设备端的数据传输来隐藏数据传输。要实现使用其余操做隐藏数据传输须要使用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运行上面示例代码的时间线：

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

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

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

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支持的每一个线程单独的默认流。尤为是在写库函数时，这一点尤其重要。若是在库函数中使用默认流，那么对于库函数用户就不会有机会实现数据传输和核函数执行的重叠了。

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