[譯]CUDA C/C++如何優化數據傳輸

本文翻譯自NVIDIA官方博客Parallel Forall,內容僅供參考,若有疑問請訪問原網站:https://devblogs.nvidia.com/p...html

文章中,咱們已經爲如何優化CUDA C/C++代碼系列文章的主要內容作了鋪墊。在這篇和以後的文章中,咱們會討論如何在提升主機和設備之間數據傳輸效率方面進行代碼優化。設備內存和GPU之間的最大帶寬(例如NVIDIA Tesla C2050的帶寬爲144 GB/s)遠大於主機內存和設備內存(如PCIe x16 Gen2總線的帶寬爲8 GB/s)之間的最大帶寬。
這個差別就意味着主機和設備之間的數據傳輸速度將成爲程序總體性能的主要瓶頸。首先讓咱們來看一看主機設備數據傳輸的一些通用準則。linux

譯者注:這裏說到的三篇文章,分別是cuda的介紹和入門、如何衡量代碼性能以及如何獲取設備的相關屬性和錯誤處理。第一篇主要是一些入門的東西,比較簡單,你們能夠自行了解;第二篇筆者的專欄已有譯文;第三篇的內容也相對簡單,若是筆者有精力的話,也會爲你們翻譯或者整理出來。git

  • 儘可能減小主機和設備之間數據的傳輸量,即便相比在CPU上,GPU上的核函數提速不多或者沒有提速也沒有關係。github

  • 使用頁鎖定主機內存(也叫作固定內存)能夠得到更高的數據傳輸帶寬。編程

  • 將多個小的數據傳輸合併爲一次更大的數據傳輸,由於這樣能夠消除每次傳輸的大部分開銷。segmentfault

  • 主機設備之間的數據傳輸有時能夠被核函數執行或者其餘數據傳輸隱藏。windows

在這篇文章咱們主要研究前三個準則,最後一個隱藏數據傳輸將會在下一篇中討論。首先咱們來討論一下如何在不修改源碼的狀況下,測量出數據傳輸的時間。api

使用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來編譯,而後以程序的名字爲參數運行nvprofapp

$ 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

譯者注:固定主機內存(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事件或者其餘計時器函數。

這篇文章的重點是如何高效地進行數據傳輸。下一篇文章,咱們會討論如何用計算和其餘數據傳輸來隱藏數據傳輸。

相關文章
相關標籤/搜索