主機(jī)和設(shè)備之間的傳輸是 GPU 計(jì)算中數(shù)據(jù)移動(dòng)最慢的一個(gè)環(huán)節(jié)，所以您應(yīng)該注意盡量減少傳輸。遵循這篇文章中的指導(dǎo)方針可以幫助你確保必要的轉(zhuǎn)移是有效的。當(dāng)您移植或編寫新的 CUDA C / C ++代碼時(shí)，我建議您從現(xiàn)有主機(jī)指針開始可分頁的傳輸。正如我前面提到的，當(dāng)您編寫更多的設(shè)備代碼時(shí)，您將消除一些中間傳輸，因此您在移植早期所花費(fèi)的優(yōu)化傳輸?shù)娜魏闻Χ伎赡鼙焕速M(fèi)。另外，我建議您不要使用 CUDA 事件或其他計(jì)時(shí)器插入代碼來測(cè)量每次傳輸所花費(fèi)的時(shí)間，而是建議您使用 nvprof， 命令行 CUDA 探查器，或者使用可視化分析工具，如 NVIDIA 可視化探查器（也包括在 CUDA 工具箱中）。

這篇文章的重點(diǎn)是提高數(shù)據(jù)傳輸?shù)男省Ｔ?下一篇文章 中，我們討論了如何將數(shù)據(jù)傳輸與計(jì)算和其他數(shù)據(jù)傳輸重疊。

在 C + C ++系列 之前的 帖子 中，我們?yōu)樵撓盗械闹饕屏Φ於嘶A(chǔ)：如何優(yōu)化 CUDA C / C ++代碼。本文就如何在主機(jī)和主機(jī)之間高效地傳輸數(shù)據(jù)展開討論。設(shè)備內(nèi)存和 GPU 之間的峰值帶寬遠(yuǎn)高于主機(jī)內(nèi)存和設(shè)備內(nèi)存之間的峰值帶寬（例如，在 GPU NVIDIA C2050 上為 144 GB / s ），而在 PCIe x16 Gen2 上為 8 GB / s 。這種差異意味著主機(jī)和 GPU 設(shè)備之間的數(shù)據(jù)傳輸?shù)膶?shí)現(xiàn)可能會(huì)影響或破壞應(yīng)用程序的整體性能。讓我們從主機(jī)數(shù)據(jù)傳輸?shù)囊话阍瓌t開始。

盡可能減少主機(jī)和設(shè)備之間傳輸?shù)臄?shù)據(jù)量，即使這意味著在 GPU 上運(yùn)行內(nèi)核，與在主機(jī) CPU 上運(yùn)行內(nèi)核相比，其速度幾乎沒有或幾乎沒有。

使用頁鎖定（或“固定”）內(nèi)存時(shí)，主機(jī)和設(shè)備之間的帶寬可能更高。

將許多小的傳輸批處理到一個(gè)較大的傳輸中執(zhí)行得更好，因?yàn)樗嗣總€(gè)傳輸?shù)拇蟛糠珠_銷。

主機(jī)和設(shè)備之間的數(shù)據(jù)傳輸有時(shí)可能與內(nèi)核執(zhí)行和其他數(shù)據(jù)傳輸重疊。

在這篇文章中，我們將研究上面的前三條準(zhǔn)則，并在下一篇文章中專門討論重疊數(shù)據(jù)傳輸。首先，我想談?wù)勅绾卧诓恍薷脑创a的情況下測(cè)量數(shù)據(jù)傳輸所花費(fèi)的時(shí)間。

用 nvprof 測(cè)量數(shù)據(jù)傳輸時(shí)間

為了測(cè)量每次數(shù)據(jù)傳輸所花費(fèi)的時(shí)間，我們可以在每次傳輸前后記錄一個(gè) 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編譯它，然后用程序文件名作為參數(shù)運(yùn)行nvprof。

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

當(dāng)我在臺(tái)式電腦上運(yùn)行時(shí)，它有一個(gè) 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 測(cè)量每個(gè) CUDA memcpy 調(diào)用所花費(fèi)的時(shí)間。它報(bào)告每個(gè)調(diào)用的平均、最小和最長時(shí)間（因?yàn)槲覀冎贿\(yùn)行每個(gè)副本一次，所有時(shí)間都是相同的）。 nvprof 非常靈活，所以請(qǐng)確保 查看文檔 。

nvprof 是 CUDA 5 中的新功能。如果您使用的是早期版本的 CUDA ，那么可以使用舊的“命令行分析器”，正如 Greg Ruetsch 在他的文章 如何在 CUDA Fortran 中優(yōu)化數(shù)據(jù)傳輸 中所解釋的那樣。

最小化數(shù)據(jù)傳輸

我們不應(yīng)該只使用內(nèi)核的 GPU 執(zhí)行時(shí)間相對(duì)于其 CPU 實(shí)現(xiàn)的執(zhí)行時(shí)間來決定是運(yùn)行 GPU 還是 CPU 版本。我們還需要考慮在 PCI-e 總線上移動(dòng)數(shù)據(jù)的成本，尤其是當(dāng)我們最初將代碼移植到 CUDA 時(shí)。因?yàn)?CUDA 的異構(gòu)編程模型同時(shí)使用了 CPU 和 GPU ，代碼可以一次移植到 CUDA 一個(gè)內(nèi)核。在移植的初始階段，數(shù)據(jù)傳輸可能支配整個(gè)執(zhí)行時(shí)間。將數(shù)據(jù)傳輸所花費(fèi)的時(shí)間與內(nèi)核執(zhí)行的時(shí)間分開記錄是值得的。正如我們已經(jīng)演示過的，使用命令行探查器很容易做到這一點(diǎn)。隨著我們移植更多的代碼，我們將刪除中間傳輸并相應(yīng)地減少總體執(zhí)行時(shí)間。

固定主機(jī)內(nèi)存

默認(rèn)情況下，主機(jī)（ CPU ）的數(shù)據(jù)分配是可分頁的。 GPU 無法直接從可分頁主機(jī)內(nèi)存訪問數(shù)據(jù)，因此當(dāng)調(diào)用從可分頁主機(jī)內(nèi)存到設(shè)備內(nèi)存的數(shù)據(jù)傳輸時(shí)， CUDA 驅(qū)動(dòng)程序必須首先分配一個(gè)臨時(shí)頁鎖定或“固定”主機(jī)數(shù)組，將主機(jī)數(shù)據(jù)復(fù)制到固定數(shù)組，然后將數(shù)據(jù)從固定數(shù)組傳輸?shù)皆O(shè)備內(nèi)存，如下圖所示。

如圖中所示，固定內(nèi)存用作從設(shè)備到主機(jī)的傳輸?shù)呐R時(shí)區(qū)域。通過直接將主機(jī)數(shù)組分配到固定內(nèi)存中，可以避免在可分頁主機(jī)數(shù)組和固定主機(jī)數(shù)組之間進(jìn)行傳輸?shù)拈_銷。使用 CUDA 或 cudaHostAlloc（） 在 CUDA C / C ++中分配被鎖定的主機(jī)內(nèi)存，并用 cudaFreeHost（） 解除它。固定內(nèi)存分配可能會(huì)失敗，因此應(yīng)該始終檢查錯(cuò)誤。下面的代碼摘要演示如何分配固定內(nèi)存以及錯(cuò)誤檢查。

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

使用主機(jī)固定內(nèi)存的數(shù)據(jù)傳輸使用與可分頁內(nèi)存?zhèn)鬏斚嗤?a target="_blank">cudaMemcpy()語法。我們可以使用下面的“帶寬測(cè)試”程序（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;
}

數(shù)據(jù)傳輸速率取決于主機(jī)系統(tǒng)的類型（主板， CPU 和芯片組）以及 GPU 。在我的筆記本電腦上，它有 Intel Core i7-2620MCPU （ 2 . 7GHz ， 2 個(gè) Sandy Bridge 內(nèi)核， 4MB L3 緩存）和 NVIDIA NVS 4200MGPU （ 1 費(fèi)米 SM ，計(jì)算能力 2 . 1 ， PCI-e Gen2 x16 ），運(yùn)行BandwidthTest會(huì)產(chǎn)生以下結(jié)果。如您所見，固定傳輸?shù)乃俣仁强煞猪搨鬏數(shù)膬杀抖唷?/p>

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 處理器。這大概是因?yàn)楦斓?CPU （和芯片組）降低了主機(jī)端的內(nèi)存復(fù)制成本。

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

不應(yīng)過度分配固定內(nèi)存。這樣做會(huì)降低整體系統(tǒng)性能，因?yàn)檫@會(huì)減少操作系統(tǒng)和其他程序可用的物理內(nèi)存量。多少是太多是很難預(yù)先判斷的，所以對(duì)于所有優(yōu)化，測(cè)試您的應(yīng)用程序和它們運(yùn)行的系統(tǒng)，以獲得最佳性能參數(shù)。

批量小轉(zhuǎn)移

由于與每個(gè)傳輸相關(guān)聯(lián)的開銷，最好將多個(gè)小傳輸一起批處理到單個(gè)傳輸中。通過使用一個(gè)臨時(shí)數(shù)組（最好是固定的）并將其與要傳輸?shù)臄?shù)據(jù)打包，這很容易做到。

對(duì)于二維數(shù)組傳輸，可以使用 cudaMemcpy2D（） 。

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

這里的參數(shù)是指向第一個(gè)目標(biāo)元素和目標(biāo)數(shù)組間距的指針，指向第一個(gè)源元素和源數(shù)組間距的指針，要傳輸?shù)淖泳仃嚨膶挾群透叨?，以?memcpy 類型。還有一個(gè) cudaMemcpy3D（） 函數(shù)用于傳輸秩為三的數(shù)組部分。

關(guān)于作者

Mark Harris 是 NVIDIA 杰出的工程師，致力于 RAPIDS 。 Mark 擁有超過 20 年的 GPUs 軟件開發(fā)經(jīng)驗(yàn)，從圖形和游戲到基于物理的模擬，到并行算法和高性能計(jì)算。當(dāng)他還是北卡羅來納大學(xué)的博士生時(shí)，他意識(shí)到了一種新生的趨勢(shì)，并為此創(chuàng)造了一個(gè)名字： GPGPU （圖形處理單元上的通用計(jì)算）。

審核編輯：郭婷