CUDA - 如何优化CUDA C/C++中的数据传输

pla66

于 2023-06-08 23:17:17 发布

阅读量573

点赞数 1

分类专栏： CUDA 文章标签： c++ c语言

原文链接：https://developer.nvidia.com/blog/how-optimize-data-transfers-cuda-cc/

版权

CUDA 专栏收录该内容

8 篇文章 2 订阅

订阅专栏

原文链接：How to Optimize Data Transfers in CUDA C/C++

文章目录

在这个CUDA C/C++系列的前三篇文章（第一篇，第二篇，第三篇）中，我们为本系列的主旨奠定了基础：如何优化CUDA C/C++代码。在这篇文章和后续的文章中，我们从如何在主机和设备之间有效地传输数据这一方面来讨论如何优化代码。设备内存和GPU之间的峰值带宽（例如，NVIDIA Tesla C2050上为144 GB/s）远高于主机内存和设备内存之间的峰值带宽（PCIe x16 Gen2上为8 GB/s）。这种差异意味着您在主机和GPU设备之间的数据传输可能会影响整体应用程序的性能。让我们从一些关于主机-设备数据传输的通用指导方针开始。

尽可能减少主机和设备之间传输的数据量，即使这意味着在GPU上运行内核与在主机CPU上运行内核相比几乎没有加速。
当使用锁页（page-locked）（或“固定（pinned）”）内存时，主机和设备之间可能会有更高的带宽。
将许多小的传输整合成一个较大的传输会使性能好得多，因为它消除了每次传输的大部分开销。
主机和设备之间的数据传输有时可以与内核执行和其他数据传输重叠。

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

用nvprof测量数据传输时间

为了测量每次数据传输所花费的时间，我们可以在每次传输前后记录一个CUDA事件，并使用cudaEventElapsedTime()，正如我们在前一篇文章中所描述的那样。然而，我们可以通过使用nvprof（CUDA Toolkit中的命令行CUDA 分析工具，从CUDA 5开始），在不使用CUDA事件的情况下获得传输时间。让我们用下面的代码示例来尝试一下，您可以在本文的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（GK104 GPU，类似于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的一个内核。在移植的初始阶段，数据传输可能会是整个执行时间的大头。将花在数据传输上的时间与花在内核执行上的时间分开进行监视是值得的。正如我们已经演示的那样，使用命令行分析工具很容易做到这一点。随着我们移植更多的代码，我们将删除中间传输并相应地减少总体执行时间。

固定主机内存（Pinned Host Memory）

默认情况下，主机（CPU）数据分配是可分页的（pageable）。GPU无法直接从可分页主机内存访问数据，因此当发起从可分页主机内存到设备内存的数据传输时，CUDA驱动程序必须首先分配一个临时锁页（“固定”）的主机数组，将主机数据复制到固定数组，然后将数据从固定数组传输到设备内存，如下所示。
在这里插入图片描述
如图所示，固定内存用作从主机传输到设备（原文为from the device to the host，可能为作者笔误）的临时区域。我们可以通过在固定内存中直接分配主机数组来避免可分页数组和固定主机数组之间的传输成本。使用cudaMallocHost()或cudaHostAlloc()在CUDA C/C++中分配固定主机内存，并使用cudaFreeHost()释放。固定内存分配可能会失败，因此您应该始终检查错误。下面的代码演示了在分配固定内存时检查错误。

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

使用主机固定内存的数据传输与使用可分页内存的传输一样使用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: %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-2620M CPU（2.7GHz，2 Sandy Bridge cores，4MB L3 Cache）和一个NVIDIA 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）和NVIDIA GeForce GTX 680 GPU（8个Kepler SM，计算能力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)

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

总结

主机和设备之间的传输是GPU计算中数据移动最慢的环节，因此您应该注意将传输最小化。遵循这篇文章中的指导可以帮助你确保必要的转移是高效的。当您移植或编写新的CUDA C/C++代码时，我建议您从现有主机指针的可分页传输开始。正如我前面提到的，随着编写的设备代码越来越多，一些中间传输将被去除，因此您在移植早期优化传输的任何努力都可能被浪费。此外，相比于使用CUDA事件或其他计时器来测量每次传输所花费的时间，我建议您使用nvprof或其他可视化分析工具（如NVIDIA Visual Profiler，也包含在CUDA Toolkit中）。

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