如何在主机和主机之间实现数据传输优化

　　主机和设备之间的传输是 GPU 计算中数据移动最慢的一个环节，所以您应该注意尽量减少传输。遵循这篇文章中的指导方针可以帮助你确保必要的转移是有效的。当您移植或编写新的 CUDA C / C ++代码时，我建议您从现有主机指针开始可分页的传输。正如我前面提到的，当您编写更多的设备代码时，您将消除一些中间传输，因此您在移植早期所花费的优化传输的任何努力都可能被浪费。另外，我建议您不要使用 CUDA 事件或其他计时器插入代码来测量每次传输所花费的时间，而是建议您使用 nvprof， 命令行 CUDA 探查器，或者使用可视化分析工具，如 NVIDIA 可视化探查器（也包括在 CUDA 工具箱中）。

　　这篇文章的重点是提高数据传输的效率。在 下一篇文章 中，我们讨论了如何将数据传输与计算和其他数据传输重叠。

　　在 C + C ++系列 之前的 帖子 中，我们为该系列的主要推力奠定了基础：如何优化 CUDA C / C ++代码。本文就如何在主机和主机之间高效地传输数据展开讨论。设备内存和 GPU 之间的峰值带宽远高于主机内存和设备内存之间的峰值带宽（例如，在 GPU NVIDIA C2050 上为 144 GB / s ），而在 PCIe x16 Gen2 上为 8 GB / s 。这种差异意味着主机和 GPU 设备之间的数据传输的实现可能会影响或破坏应用程序的整体性能。让我们从主机数据传输的一般原则开始。

　　尽可能减少主机和设备之间传输的数据量，即使这意味着在 GPU 上运行内核，与在主机 CPU 上运行内核相比，其速度几乎没有或几乎没有。

　　使用页锁定（或“固定”）内存时，主机和设备之间的带宽可能更高。

　　将许多小的传输批处理到一个较大的传输中执行得更好，因为它消除了每个传输的大部分开销。

　　主机和设备之间的数据传输有时可能与内核执行和其他数据传输重叠。

　　在这篇文章中，我们将研究上面的前三条准则，并在下一篇文章中专门讨论重叠数据传输。首先，我想谈谈如何在不修改源代码的情况下测量数据传输所花费的时间。

　　用 nvprof 测量数据传输时间

　　为了测量每次数据传输所花费的时间，我们可以在每次传输前后记录一个 CUDA 事件，并使用 cudaEventElapsedTime（） ，正如我们所描述的 在上一篇文章中 ， CUDA 工具箱中包含的命令行 CUDA 探查器（从 CUDA 5 开始）。让我们用下面的代码示例来尝试一下，您可以在 CUDA 中找到它。

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

当我在台式电脑上运行时，它有一个 geforcegtx680 （ GK104GPU ，类似于 Tesla K10 ），我得到以下输出。

$ 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]

　　如您所见， nvprof 测量每个 CUDA memcpy 调用所花费的时间。它报告每个调用的平均、最小和最长时间（因为我们只运行每个副本一次，所有时间都是相同的）。 nvprof 非常灵活，所以请确保 查看文档 。

　　nvprof 是 CUDA 5 中的新功能。如果您使用的是早期版本的 CUDA ，那么可以使用旧的“命令行分析器”，正如 Greg Ruetsch 在他的文章 如何在 CUDA Fortran 中优化数据传输 中所解释的那样。

　　最小化数据传输

　　我们不应该只使用内核的 GPU 执行时间相对于其 CPU 实现的执行时间来决定是运行 GPU 还是 CPU 版本。我们还需要考虑在 PCI-e 总线上移动数据的成本，尤其是当我们最初将代码移植到 CUDA 时。因为 CUDA 的异构编程模型同时使用了 CPU 和 GPU ，代码可以一次移植到 CUDA 一个内核。在移植的初始阶段，数据传输可能支配整个执行时间。将数据传输所花费的时间与内核执行的时间分开记录是值得的。正如我们已经演示过的，使用命令行探查器很容易做到这一点。随着我们移植更多的代码，我们将删除中间传输并相应地减少总体执行时间。

　　固定主机内存

　　默认情况下，主机（ CPU ）的数据分配是可分页的。 GPU 无法直接从可分页主机内存访问数据，因此当调用从可分页主机内存到设备内存的数据传输时， CUDA 驱动程序必须首先分配一个临时页锁定或“固定”主机数组，将主机数据复制到固定数组，然后将数据从固定数组传输到设备内存，如下图所示。

　　如图中所示，固定内存用作从设备到主机的传输的临时区域。通过直接将主机数组分配到固定内存中，可以避免在可分页主机数组和固定主机数组之间进行传输的开销。使用 CUDA 或 cudaHostAlloc（） 在 CUDA C / C ++中分配被锁定的主机内存，并用 cudaFreeHost（） 解除它。固定内存分配可能会失败，因此应该始终检查错误。下面的代码摘要演示如何分配固定内存以及错误检查。

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

使用主机固定内存的数据传输使用与可分页内存传输相同的 cudaMemcpy() 语法。我们可以使用下面的“带宽测试”程序（ Github 上也有 ）来比较可分页和固定的传输速率。

#include 
#include 

// 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: %s\n",
            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 transfers\n", 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): %f\n", 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): %f\n", bytes * 1e-6 / time);

  for (int i = 0; i < n; ++i) {
    if (h_a[i] != h_b[i]) {
      printf("*** %s transfers failed ***\n", 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: %s\n", prop.name);
  printf("Transfer size (MB): %d\n", 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-2620MCPU （ 2 . 7GHz ， 2 个 Sandy Bridge 内核， 4MB L3 缓存）和 NVIDIA NVS 4200MGPU （ 1 费米 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

更快速的 3GHz 处理器（ 3GHz ， 3GHz ）和 3K 处理器（ 3GHz ）相比，我们可以更快地使用 3K 处理器（ 3GHz ）和 3GHz 处理器。这大概是因为更快的 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)

　　这里的参数是指向第一个目标元素和目标数组间距的指针，指向第一个源元素和源数组间距的指针，要传输的子矩阵的宽度和高度，以及 memcpy 类型。还有一个 cudaMemcpy3D（） 函数用于传输秩为三的数组部分。

　　关于作者

　　Mark Harris 是 NVIDIA 杰出的工程师，致力于 RAPIDS 。 Mark 拥有超过 20 年的 GPUs 软件开发经验，从图形和游戏到基于物理的模拟，到并行算法和高性能计算。当他还是北卡罗来纳大学的博士生时，他意识到了一种新生的趋势，并为此创造了一个名字： GPGPU （图形处理单元上的通用计算）。

　　审核编辑：郭婷