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

在前三篇文章中，我们已经为如何优化CUDA C/C++代码系列文章的主要内容做了铺垫。在这篇和之后的文章中，我们会讨论如何在提高主机和设备之间数据传输效率方面进行代码优化。设备内存和GPU之间的最大带宽(例如NVIDIA Tesla C2050的带宽为144 GB/s)远大于主机内存和设备内存(如PCIe x16 Gen2总线的带宽为8 GB/s)之间的最大带宽。
这个差异就意味着主机和设备之间的数据传输速度将成为程序整体性能的主要瓶颈。首先让我们来看一看主机设备数据传输的一些通用准则。

译者注：这里说到的三篇文章，分别是cuda的介绍和入门、如何衡量代码性能以及如何获取设备的相关属性和错误处理。第一篇主要是一些入门的东西，比较简单，大家可以自行了解；第二篇笔者的专栏已有译文；第三篇的内容也相对简单，如果笔者有精力的话，也会为大家翻译或者整理出来。

• 尽量减少主机和设备之间数据的传输量，即使相比在CPU上，GPU上的核函数提速很少或者没有提速也没有关系。
• 使用页锁定主机内存(也叫做固定内存)可以获得更高的数据传输带宽。
• 将多个小的数据传输合并为一次更大的数据传输，因为这样可以消除每次传输的大部分开销。
• 主机设备之间的数据传输有时可以被核函数执行或者其他数据传输隐藏。

在这篇文章我们主要研究前三个准则，最后一个隐藏数据传输将会在下一篇中讨论。首先我们来讨论一下如何在不修改源码的情况下，测量出数据传输的时间。

使用nvprof测量数据传输时间

正如我们在前面文章中所说的那样，我们可以在数据传输的前后使用CUDA事件记录然后使用cudaEventElapsedTime()来计算出传输的时间。其实，借助于nvprof，我们可以不需要使用CUDA事件而修改源代码就可以获取到所消耗的传输时间。这是一个命令行的CUDA分析器，CUDA 5及以后版本的CUDA toolkit都含有该软件。我们可以尝试使用一下这个软件，下面是我们的测试代码，源码可以在这篇文章的Github仓库中找到。

int main()
{
    const unsigned int N = 1048576;
    const unsigned int bytes = N * sizeof(int);
    int *h_a = (int*)malloc(bytes);
    int *d_a;
    cudaMalloc((int**)&d_a, bytes);

    memset(h_a, 0, bytes);
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(h_a, d_a, bytes, cudaMemcpyDeviceToHost);

    return 0;
}

为了可以分析该代码，我们仅仅使用nvcc来编译，然后以程序的名字为参数运行nvprof。

$ nvcc profile.cu -o profile_test
$ nvprof ./profile_test

下面是我在GeForce GTX 680显卡的电脑上运行之后的输出：

$ nvprof ./a.out
======== NVPROF is profiling a.out...
======== Command: a.out
======== Profiling result:
Time(%)     Time  Calls      Avg      Min      Max Name
  50.08 718.11us      1 718.11us 718.11us 718.11us [CUDA memcpy DtoH]
  49.92 715.94us      1 715.94us 715.94us 715.94us [CUDA memcpy HtoD]

译者注:很明显，原文作者是在linux下进行的测试。而在windows下可能会出问题，如果有问题的话，可以在上面的代码中的main函数最后加上cudaThreadExit()函数。另外，windows下可以使用图形界面的CUDA分析器——NVIDIA Visual Profiler，后面会提到。

正如你所看到的，nvprof测得了每一次内存拷贝所用的时间。它报告了每次调用的平均时间、最小时间和最大时间(因为每次拷贝我们只运行了一次，所以所有的时间都是相同的)。nvprof使用起来是相当灵活的，所以请务必查看相关文档来学习。

nvprof是CUDA 5中新添加的。所以正如Greg Ruetsch在这篇文章How to Optimize Data Transfers in CUDA Fortra中解释的，如果你使用的CUDA是更早版本的，你可以使用旧版本的“命令行分析器”。

尽量减少数据传输

我们不应该仅仅通过比较核函数在GPU上的执行时间和在CPU上的执行时间来决定是使用GPU版本还是CPU版本。我们也需要考虑数据在PCIe总线上传输的时间开销，尤其是在我们开始将代码移植到CUDA上的时候。由于CUDA异构编程模型同时使用CPU和GPU，所以代码可以一次移植到CUDA的一个核函数中。在移植的开始阶段，数据传输可能会在总体的执行时间中占主导地位。因此，我们需要关注单独的数据传输的时间。正如我们前面所演示的，使用命令行分析器可以很容易地得到这个数据。当我们移植更多的代码时，我们就会去掉中间的传输从而相应的减少总体的执行时间。

译者注：这里所说的“移植”(port)的意思是翻译，即从一种编程语言转为另一种语言，因此这里的意思就是将CPU代码转化为GPU代码。

固定主机内存

主机(CPU)数据分配的内存默认是可分页的。GPU不能直接访问可分页的主机内存，所以当从可分页内存到设备内存的进行数据传输时，CUDA驱动必须首先分配一个临时的不可分页的或者固定的主机数组，然后将主机数据拷贝到固定数组里，最后再将数据从固定数组转移到设备内存，如下图所示：

使用固定主机内存

译者注：固定主机内存(Pinned Host Memory)又称为页锁定主机内存(page-locked host memory)或者不可分页主机内存，它有一个重要属性：操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中。

正如你在图中所看到的那样，固定内存被用作数据传输的暂存区。我们可以通过直接分配固定内存的主机数组来避免这一开销。在CUDA C/C++中，我们可以使用cudaMallocHost()或者cudaHostAlloc()来分配固定内存，使用cudaFreeHost()来释放内存。固定内存的分配有可能会失败，所以你应该总是检查错误。下面的代码片段演示了如何分配固定内存并进行错误检查。

cudaError_t status = cudaMallocHost((void**)&h_aPinned, bytes);
if (status != cudaSuccess)
  printf("Error allocating pinned host memoryn");

固定内存的数据传输和可分页内存一样，使用相同的cudaMemcpy()语法。我们可以使用下面的“bandwidthtest”(带宽测试)程序(同样可以在Github上找到)来对比可分页内存和固定内存的传输速度。

#include <stdio.h>
#include <assert.h>

// Convenience function for checking CUDA runtime API results
// can be wrapped around any runtime API call. No-op in release builds.
inline
cudaError_t checkCuda(cudaError_t result)
{
#if defined(DEBUG) || defined(_DEBUG)
  if (result != cudaSuccess) {
    fprintf(stderr, "CUDA Runtime Error: %sn",
            cudaGetErrorString(result));
    assert(result == cudaSuccess);
  }
#endif
  return result;
}

void profileCopies(float        *h_a,
                   float        *h_b,
                   float        *d,
                   unsigned int  n,
                   char         *desc)
{
  printf("n%s transfersn", desc);

  unsigned int bytes = n * sizeof(float);

  // events for timing
  cudaEvent_t startEvent, stopEvent;

  checkCuda( cudaEventCreate(&startEvent) );
  checkCuda( cudaEventCreate(&stopEvent) );

  checkCuda( cudaEventRecord(startEvent, 0) );
  checkCuda( cudaMemcpy(d, h_a, bytes, cudaMemcpyHostToDevice) );
  checkCuda( cudaEventRecord(stopEvent, 0) );
  checkCuda( cudaEventSynchronize(stopEvent) );

  float time;
  checkCuda( cudaEventElapsedTime(&time, startEvent, stopEvent) );
  printf("  Host to Device bandwidth (GB/s): %fn", bytes * 1e-6 / time);

  checkCuda( cudaEventRecord(startEvent, 0) );
  checkCuda( cudaMemcpy(h_b, d, bytes, cudaMemcpyDeviceToHost) );
  checkCuda( cudaEventRecord(stopEvent, 0) );
  checkCuda( cudaEventSynchronize(stopEvent) );

  checkCuda( cudaEventElapsedTime(&time, startEvent, stopEvent) );
  printf("  Device to Host bandwidth (GB/s): %fn", bytes * 1e-6 / time);

  for (int i = 0; i < n; ++i) {
    if (h_a[i] != h_b[i]) {
      printf("*** %s transfers failed ***", desc);
      break;
    }
  }

  // clean up events
  checkCuda( cudaEventDestroy(startEvent) );
  checkCuda( cudaEventDestroy(stopEvent) );
}

int main()
{
  unsigned int nElements = 4*1024*1024;
  const unsigned int bytes = nElements * sizeof(float);

  // host arrays
  float *h_aPageable, *h_bPageable;   
  float *h_aPinned, *h_bPinned;

  // device array
  float *d_a;

  // allocate and initialize
  h_aPageable = (float*)malloc(bytes);                    // host pageable
  h_bPageable = (float*)malloc(bytes);                    // host pageable
  checkCuda( cudaMallocHost((void**)&h_aPinned, bytes) ); // host pinned
  checkCuda( cudaMallocHost((void**)&h_bPinned, bytes) ); // host pinned
  checkCuda( cudaMalloc((void**)&d_a, bytes) );           // device

  for (int i = 0; i < nElements; ++i) h_aPageable[i] = i;      
  memcpy(h_aPinned, h_aPageable, bytes);
  memset(h_bPageable, 0, bytes);
  memset(h_bPinned, 0, bytes);

  // output device info and transfer size
  cudaDeviceProp prop;
  checkCuda( cudaGetDeviceProperties(&prop, 0) );

  printf("nDevice: %sn", prop.name);
  printf("Transfer size (MB): %dn", bytes / (1024 * 1024));

  // perform copies and report bandwidth
  profileCopies(h_aPageable, h_bPageable, d_a, nElements, "Pageable");
  profileCopies(h_aPinned, h_bPinned, d_a, nElements, "Pinned");

  printf("n");

  // cleanup
  cudaFree(d_a);
  cudaFreeHost(h_aPinned);
  cudaFreeHost(h_bPinned);
  free(h_aPageable);
  free(h_bPageable);

  return 0;
}

数据传输速度可能会取决于不同的主机系统(主板、CPU和芯片组)以及GPU。在我的笔记本电脑(Intel Core i7-2620M CPU (2.7GHz, 2 Sandy Bridge cores, 4MB L3 Cache)和一个英伟达NVS 4200M GPU (1 Fermi SM,计算能力2.1, PCI-e Gen2 x16))上，BandwidthTest运行的结果如下，正如你所看到的固定内存的传输速度是可分页内存的两倍多。

Device: NVS 4200M
Transfer size (MB): 16

Pageable transfers
  Host to Device bandwidth (GB/s): 2.308439
  Device to Host bandwidth (GB/s): 2.316220

Pinned transfers
  Host to Device bandwidth (GB/s): 5.774224
  Device to Host bandwidth (GB/s): 5.958834

在我的台式电脑上(Intel Core i7-3930K CPU (3.2 GHz, 6 Sandy Bridge cores, 12MB L3 Cache)和1个NVIDIA GeForce GTX 680 GPU (8 Kepler SMs,计算能力3.0))，可分页内存的数据传输就相当快了，下面是输出的结果。这很可能是因为高速的CPU(和芯片组)减少了主机端内存拷贝的开销。

Device: GeForce GTX 680
Transfer size (MB): 16

Pageable transfers
  Host to Device bandwidth (GB/s): 5.368503
  Device to Host bandwidth (GB/s): 5.627219

Pinned transfers
  Host to Device bandwidth (GB/s): 6.186581
  Device to Host bandwidth (GB/s): 6.670246

你应该避免分配过多的固定内存。分配过多的固定内存会降低系统的整体性能，因为这会减少操作系统和其他程序可用的物理内存空间。我们很难知道到底多少才算太多，所以和其他所有优化一样，我们需要对程序和系统进行测试才能获得最优的性能参数。

合并小规模的数据传输

因为每次数据传输都会产生额外的开销，所以最好将多个小规模的数据传输合并为单独的一次数据传输。我们可以使用临时的数组，然后用将要传输的数据填充该数组即可，而且最好使用固定内存的数组。

对于二维数组的传输，你可以使用cudaMemcpy2D()。

cudaMemcpy2D(dest, dest_pitch, src, src_pitch, w, h, cudaMemcpyHostToDevice)

这个函数的参数分别是指向目标内存第一个元素的指针、目标数组的步长(pitch)、指向源内存第一个元素的指针、源数组的步长(pitch)、要传输的子矩阵的宽和高、内存拷贝的类型。另外还有一个函数cudaMemcpy3D()可以用于三维数组段传输。

总结

主机设备间的数据传输是GPU计算中最慢的数据移动环节，所以你应该注意尽量减少它们之间的传输。采用本文的这些准则可以使你高效地进行数据传输。当你移植或者编写新的CUDA C/C++代码时，我推荐你先使用可分页的数据传输方式。正如我之前说的，当你的设备代码越来越多时，你就可以消除一些中间的数据传输，所以过早的数据传输优化很可能会白费。

译者注：正如高德纳所说：

过早的优化是万恶之源！

另外，对于数据传输的时间的测量，我推荐你使用命令行CUDA分析器——nvprof或者其他可视化分析器如英伟达的Visual Profiler(也被包含在CUDA toolkit中)，而不是使用CUDA事件或者其他计时器函数。

这篇文章的重点是如何高效地进行数据传输。下一篇文章，我们会讨论如何用计算和其他数据传输来隐藏数据传输。