[譯]在CUDA C/C++中如何隱藏數據傳輸

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

在上一篇博客中，咱們討論瞭如何在主機和設備之間高效地進行數據傳輸。在這篇文章中，咱們將討論如何使用主機端的計算、設備端的計算以及某些狀況下的主機與設備端的數據傳輸來隱藏數據傳輸。要實現使用其餘操做隱藏數據傳輸須要使用CUDA流，因此首先讓咱們來了解一下CUDA流。

譯者注：這裏爲了符合中文的習慣，我將「Overlap Data Transfers」譯爲「隱藏數據傳輸」。「overlap」，原意爲重疊，這裏將其翻譯爲隱藏，既能夠表達隱藏了數據傳輸的開銷，也能夠隱含地表達重疊的意思，更加的形象貼切。可是某些地方，爲了表達順暢，我也將其直接翻譯爲重疊。無論翻譯成什麼，只須要明白隱藏就是靠重疊來實現的，經過將幾種相同或不一樣的操做重疊，咱們就能夠近似地實現隱藏某些開銷。

CUDA流

CUDA流是由主機端發佈，在設備端順序執行的一系列操做。在一個CUDA流中的操做能夠保證按既定的順序執行，而在不一樣的流中的操做能夠交疊執行，有時甚至能夠併發(concurrently)執行。

默認流

全部設備操做，包括核函數和數據傳輸，都運行在CUDA流中。當沒有指定使用哪一個流時，就會使用默認流(也叫作「空流」，null stream)。默認流不一樣於其餘流，由於它是一個對於設備上操做同步的CUDA流：直到以前發佈在流中的全部操做完成，默認流中的操做纔會開始；默認流中的操做必須在其餘流中的操做開始前完成。

請注意在2015年發佈的CUDA 7引入了一個新的特性——能夠在每一個主機線程中使用單獨的默認流；也能夠將每一個線程的默認流做爲普通流使用(即它們不對其餘流中的操做進行同步)。詳情請閱讀這篇文章——GPU Pro Tip: CUDA 7 Streams Simplify Concurrency。

讓咱們來一塊兒看一個使用默認流的簡單例子，以及討論如何從主機和設備的角度分析流中操做的執行過程。

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a)
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

在上面的代碼中，從設備的角度來看，全部上述三個操做都被髮布在相同的流——默認流中，它們會按照發布的順序執行。從主機的角度來看，隱式的數據傳輸是同步的，而核函數啓動是異步的。既然主機到設備的數據傳輸(第一行)是同步的，那麼等到數據傳輸完成CPU線程纔會調用核函數。一旦核函數被調用，CPU線程會馬上執行到第三行，可是因爲設備端的執行順序這行的數據傳輸並不會馬上開始。

從主機的角度來看，核函數執行的異步行爲很是有利於設備和主機端的計算重疊。咱們能夠在上面的代碼中添加一些獨立的CPU計算。

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a)
myCpuFunction(b)
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

在上面的代碼中，一旦increment()核函數在設備端被調用，CPU線程就會馬上執行myCpuFunction()，這樣就實現了主機端myCpuFunction執行與設備端核函數執行的重疊。不管是主機端的函數先執行仍是設備端的核函數先執行都不會影響以後設備到主機的數據傳輸，由於只有在核函數執行完畢以後它纔會開始。從設備的角度來看，與前一個代碼相比什麼也沒有改變，設備徹底不會意識到myCpuFunction()的執行。

非默認流

非默認流在主機端聲明、建立、銷燬的C/C++代碼以下：

cudaStream_t stream1;
cudaError_t result;
result = cudaStreamCreate(&stream1)
result = cudaStreamDestroy(stream1)

咱們可使用cudaMemcpyAsync()函數來在一個非默認流中發佈一個數據傳輸，這很相似於以前博客中討論的cudaMemcpy()函數，區別就在於前者有第四個參數，用於標識使用哪一個CUDA流。

result = cudaMemcpyAsync(d_a, a, N, cudaMemcpyHostToDevice, stream1)

cudaMemcpyAsync()在主機端是非同步的，因此當數據傳輸一旦開始控制權就會馬上返回到主機線程。對於2D和3D的數組的拷貝，我麼可使用 cudaMemcpy2DAsync()和cudaMemcpy3DAsync()的函數形式。

在啓動核函數時，咱們須要使用第四個執行時配置參數(三對尖括號中)——流標識符(第三個執行時配置參數是爲了分配共享內存，咱們會在以後討論，這裏使用0)。

increment<<<1,N,0,stream1>>>(d_a)

流的同步

你可能會遇到須要將主機代碼與流中操做同步的狀況，可是非默認流中的全部操做對於主機代碼都是非同步的。有好幾種方法能夠解決這個問題。最有力的方法是使用 cudaDeviceSynchronize()，它能夠阻塞主機代碼直到以前全部發布在設備端的代碼所有完成爲止。在大多數狀況下，這其實都太過了，並且也會有損程序性能，由於這種方式會拖延整個設備和主機線程。

譯者注：流的同步通常被用於時間測量。

CUDA流API中有多種溫和的方式來同步主機代碼。函數 cudaStreamSynchronize(流)能夠用於阻塞主機線程直到以前發佈在指定流的全部操做完成爲止。函數cudaStreamQuery(流)能夠用於測試以前發佈在指定流的全部操做是否完成，但不會阻塞主機線程。函數cudaEventSynchronize(事件)和cudaEventQuery(事件)與前兩種函數很像，區別在於後者是基於指定事件是否被記錄而前者是基於指定的流是否空閒。你也能夠在一個單獨的流中基於一個特定的事件使用cudaStreamWaitEvent(事件)函數(即便事件被記錄在不一樣的流中或者不一樣的設備中！)

核函數執行和數據傳輸的重疊

以前咱們已經演示瞭如何在默認流中用主機端代碼來隱藏核函數執行。可是咱們的主要目的是演示如何用核函數執行隱藏數據傳輸。要實現它有幾點要求：

• 設備必須能夠「併發地拷貝和執行」。咱們能夠經過訪問cudaDeviceProp結構體的deviceOverlap屬性或者從CUDA SDK/Toolkit中deviceQuery示例程序的輸出中得到。幾乎全部計算能力1.1及以上的設備都支持設備重疊。

• 核函數執行和數據傳輸必須在不一樣的非默認流中。

• 涉及到數據傳輸的主機內存必須是固定主機內存。

下面讓咱們來修改上面的代碼以使用多個CUDA流，看一看是否實現了數據傳輸的隱藏。完整的代碼能夠在Github上找到。在這個被修改的代碼中，咱們將大小爲N的數組分爲streamSize大小的數據塊。既然核函數能夠獨立地操做全部數據，那麼每一個數據塊也能夠被獨立地處理。流(非默認流)的數量nStreams=N/streamSize。實現數據的分解處理有多種方式，一種是將對每一個數據塊的全部操做都放到一個循環中，代碼以下所示：

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  cudaMemcpyAsync(&d_a[offset], &a[offset], streamBytes, cudaMemcpyHostToDevice, stream[i]);
  kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
  cudaMemcpyAsync(&a[offset], &d_a[offset], streamBytes, cudaMemcpyDeviceToHost, stream[i]);
}

另外一種方式是將相似的操做放在一塊兒批處理，首先發布全部主機到設備的數據傳輸，以後是核函數執行，而後就是設備到主機的數據傳輸，代碼以下所示：

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  cudaMemcpyAsync(&d_a[offset], &a[offset],
                  streamBytes, cudaMemcpyHostToDevice, cudaMemcpyHostToDevice, stream[i]);
}

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
}

for (int i = 0; i < nStreams; ++i) {
  int offset = i * streamSize;
  cudaMemcpyAsync(&a[offset], &d_a[offset],
                  streamBytes, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToHost, stream[i]);
}

上述兩種異步方法都會產生正確的結果，並且同一個流中相互依賴的操做都會按照須要的順序執行。然而，這兩種方式的性能在不一樣版本的GPU上具備很大的差別。在Tesla C1060的GPU(計算能力1.3)上運行上述測試代碼，結果以下：

Device : Tesla C1060

Time for sequential transfer and execute (ms ): 12.92381
  max error : 2.3841858E -07
Time for asynchronous V1 transfer and execute (ms ): 13.63690
  max error : 2.3841858E -07
Time for asynchronous V2 transfer and execute (ms ): 8.84588
  max error : 2.3841858E -07

在Tesla C2050(計算能力2.0)，咱們獲得如下結果：

Device : Tesla C2050

Time for sequential transfer and execute (ms ): 9.984512
  max error : 1.1920929e -07
Time for asynchronous V1 transfer and execute (ms ): 5.735584
  max error : 1.1920929e -07
Time for asynchronous V2 transfer and execute (ms ): 7.597984
  max error : 1.1920929e -07

這裏數據傳輸和核函數順序執行的同步版本能夠做爲比較上述兩種異步版本是否有加速效果的基準。爲何這兩種異步執行策略在不一樣架構上的效果不一樣呢？爲了解釋這一結果，咱們須要瞭解CUDA設備如何調度和執行任務。CUDA設備中存在多種不一樣任務的引擎，它們會對發佈的操做進行排隊。它們的功能就是維護不一樣引擎中任務間的依賴，可是在引擎內部全部的外部依賴都會丟失；每一個引擎中的任務都會按照它們被髮布的順序執行。C1060有一個單獨的拷貝引擎和一個單獨的核函數引擎。下圖是C1060運行上面示例代碼的時間線：

NOTE:H2D表示主機到設備；D2H表示設備到主機

在這個原理圖中，咱們假設主機到設備的數據傳輸、核函數執行、設備到主機三者所用的時間相同(所選擇的核函數代碼就是專門這樣設計的)。正如預料的那樣，順序執行的核函數並無任何操做重疊。對於異步版本1的代碼，拷貝引擎中的執行順序是： H2D 1號流, D2H 1號流, H2D 2號流, D2H 2號流, 以此類推。這就是爲何異步版本1沒有任何加速的緣由：在拷貝引擎上任務的發佈順序使得核函數執行和數據傳輸沒法重疊。然而，從版本2較少的執行時間來看，全部主機到設備的數據傳輸都在設備到主機的數據傳輸以前，是有可能實現重疊的。在原理圖中，咱們能夠看出異步版本理論時間是順序版本的8/12，前面的結果8.7ms恰好符合這個推算。

在C2050中，有兩個特徵共同致使了它與C1060的性能差別。C2050有兩個拷貝引擎，一個是用於主機到設備的數據傳輸，另外一個用於設備到主機的數據傳輸，第三個引擎是核函數引擎。下圖描述了C2050執行示例代碼的時間線：

C2050具備兩個拷貝引擎剛好解釋了爲何異步版本1在C2050上具備很好的加速效果：與C1060正相反，在stream[i]上設備到主機的數據傳輸並不會妨礙stream[i+1]上的主機到設備的數據傳輸，由於在C2050上每一個方向的拷貝都有單獨的引擎。上面的原理圖顯示，該異步版本1的執行時間大約是順序版本的一半，和實際結果相差無幾。

可是咱們該如何解釋異步版本2在C2050上的性能降低呢？其實這與C2050能夠併發執行多個核函數有關。當多個核函數背靠背地被髮布在不一樣的流(非默認流)中時，調度器會盡力確保這些核函數併發執行，結果就致使每一個核函數完成的信號被延遲，即全部核函數執行完畢才發出信號，而這個信號負責啓動設備到主機的數據傳輸。所以，在異步版本2中，主機到設備的數據傳輸與核函數執行能夠重疊，而核函數執行與設備到主機的數據傳輸不能重疊。上面的原理圖中顯示異步版本2的整體時間大約是順序版本的9/12，正好與實驗結果7.5ms相吻合。

關於這個例子，在這篇文章CUDA Fortran Asynchronous Data Transfers中有更詳細的講解。讓人高興的是，對於計算能力3.5的設備(K20系列)，它所具備的超Q特性使得咱們已經不在須要特別安排啓動順序，因此上述兩個版本都會有很好的加速效果。咱們會在未來的博客中討論如何使用開普勒的這些特性。可是如今讓咱們來看一下Tesla K20c GPU的運行結果。正如你所看到的，兩個異步執行版本相比同步版本都有相同的加速效果。

Device : Tesla K20c
Time for sequential transfer and execute (ms): 7.101760
  max error : 1.1920929e -07
Time for asynchronous V1 transfer and execute (ms): 3.974144
  max error : 1.1920929e -07
Time for asynchronous V2 transfer and execute (ms): 3.967616
  max error : 1.1920929e -07

總結

這篇文章和以前的文章都對如何優化主機和設備間的數據傳輸進行了討論。以前的文章強調如何儘量減小數據傳輸等任務的執行時間，這篇文章介紹了流以及如何使用它們來隱藏數據傳輸，即併發地執行數據拷貝和核函數。

說到流，我必需要提醒一點：儘管使用默認流很是的方便並且代碼寫起來也很簡單，但咱們仍是應該使用非默認流或者CUDA 7支持的每一個線程單獨的默認流。尤爲是在寫庫函數時，這一點尤其重要。若是在庫函數中使用默認流，那麼對於庫函數用戶就不會有機會實現數據傳輸和核函數執行的重疊了。

如今你應該明白瞭如何高效地在主機和設備間傳輸數據，在下一篇博客中咱們開始學習如何在覈函數中高效的訪問數據。