本文翻譯自NVIDIA官方博客Parallel Forall，內(nèi)容僅供參考昂验，如有疑問請訪問原網(wǎng)站:https://devblogs.nvidia.com/parallelforall/how-overlap-data-transfers-cuda-cc/

在上一篇博客中戈擒，我們討論了如何在主機和設備之間高效地進行數(shù)據(jù)傳輸眶明。在這篇文章中，我們將討論如何使用主機端的計算峦甩、設備端的計算以及某些情況下的主機與設備端的數(shù)據(jù)傳輸來隱藏數(shù)據(jù)傳輸赘来。要實現(xiàn)使用其他操作隱藏數(shù)據(jù)傳輸需要使用CUDA流现喳，所以首先讓我們來了解一下CUDA流。

譯者注：這里為了符合中文的習慣犬辰，我將“Overlap Data Transfers”譯為“隱藏數(shù)據(jù)傳輸”嗦篱。“overlap”幌缝，原意為重疊灸促，這里將其翻譯為隱藏，既可以表達隱藏了數(shù)據(jù)傳輸?shù)拈_銷涵卵，也可以隱含地表達重疊的意思浴栽，更加的形象貼切。但是某些地方轿偎，為了表達順暢典鸡，我也將其直接翻譯為重疊。不管翻譯成什么坏晦，只需要明白隱藏就是靠重疊來實現(xiàn)的萝玷，通過將幾種相同或不同的操作重疊，我們就可以近似地實現(xiàn)隱藏某些開銷昆婿。

CUDA流

CUDA流是由主機端發(fā)布球碉，在設備端順序執(zhí)行的一系列操作。在一個CUDA流中的操作可以保證按既定的順序執(zhí)行仓蛆，而在不同的流中的操作可以交疊執(zhí)行睁冬，有時甚至可以并發(fā)(concurrently)執(zhí)行。

默認流

所有設備操作看疙，包括核函數(shù)和數(shù)據(jù)傳輸豆拨，都運行在CUDA流中。當沒有指定使用哪個流時狼荞，就會使用默認流(也叫做“空流”辽装，null stream)帮碰。默認流不同于其他流相味，因為它是一個對于設備上操作同步的CUDA流：直到之前發(fā)布在流中的所有操作完成，默認流中的操作才會開始殉挽；默認流中的操作必須在其他流中的操作開始前完成丰涉。

請注意在2015年發(fā)布的CUDA 7引入了一個新的特性——可以在每個主機線程中使用單獨的默認流；也可以將每個線程的默認流作為普通流使用(即它們不對其他流中的操作進行同步)斯碌。詳情請閱讀這篇文章——GPU Pro Tip: CUDA 7 Streams Simplify Concurrency一死。

讓我們來一起看一個使用默認流的簡單例子，以及討論如何從主機和設備的角度分析流中操作的執(zhí)行過程傻唾。

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

在上面的代碼中投慈，從設備的角度來看承耿，所有上述三個操作都被發(fā)布在相同的流——默認流中，它們會按照發(fā)布的順序執(zhí)行伪煤。從主機的角度來看加袋，隱式的數(shù)據(jù)傳輸是同步的，而核函數(shù)啟動是異步的抱既。既然主機到設備的數(shù)據(jù)傳輸(第一行)是同步的职烧，那么等到數(shù)據(jù)傳輸完成CPU線程才會調(diào)用核函數(shù)。一旦核函數(shù)被調(diào)用防泵，CPU線程會立刻執(zhí)行到第三行蚀之，但是由于設備端的執(zhí)行順序這行的數(shù)據(jù)傳輸并不會立刻開始。

從主機的角度來看捷泞，核函數(shù)執(zhí)行的異步行為非常有利于設備和主機端的計算重疊足删。我們可以在上面的代碼中添加一些獨立的CPU計算。

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

在上面的代碼中锁右，一旦increment()核函數(shù)在設備端被調(diào)用壹堰，CPU線程就會立刻執(zhí)行myCpuFunction()，這樣就實現(xiàn)了主機端myCpuFunction執(zhí)行與設備端核函數(shù)執(zhí)行的重疊骡湖。無論是主機端的函數(shù)先執(zhí)行還是設備端的核函數(shù)先執(zhí)行都不會影響之后設備到主機的數(shù)據(jù)傳輸贱纠，因為只有在核函數(shù)執(zhí)行完畢之后它才會開始。從設備的角度來看响蕴，與前一個代碼相比什么也沒有改變谆焊，設備完全不會意識到myCpuFunction()的執(zhí)行。

非默認流

非默認流在主機端聲明浦夷、創(chuàng)建辖试、銷毀的C/C++代碼如下：

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

我們可以使用cudaMemcpyAsync()函數(shù)來在一個非默認流中發(fā)布一個數(shù)據(jù)傳輸，這很類似于之前博客中討論的cudaMemcpy()函數(shù)劈狐，區(qū)別就在于前者有第四個參數(shù)罐孝，用于標識使用哪個CUDA流。

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

cudaMemcpyAsync()在主機端是非同步的肥缔，所以當數(shù)據(jù)傳輸一旦開始控制權就會立刻返回到主機線程莲兢。對于2D和3D的數(shù)組的拷貝，我么可以使用cudaMemcpy2DAsync()和cudaMemcpy3DAsync()的函數(shù)形式续膳。

在啟動核函數(shù)時改艇，我們需要使用第四個執(zhí)行時配置參數(shù)(三對尖括號中)——流標識符(第三個執(zhí)行時配置參數(shù)是為了分配共享內(nèi)存，我們會在之后討論坟岔，這里使用0)谒兄。

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

流的同步

你可能會遇到需要將主機代碼與流中操作同步的情況，但是非默認流中的所有操作對于主機代碼都是非同步的社付。有好幾種方法可以解決這個問題承疲。最有力的方法是使用cudaDeviceSynchronize()邻耕，它可以阻塞主機代碼直到之前所有發(fā)布在設備端的代碼全部完成為止。在大多數(shù)情況下燕鸽，這其實都太過了赊豌，而且也會有損程序性能，因為這種方式會拖延整個設備和主機線程绵咱。

譯者注：流的同步一般被用于時間測量碘饼。

CUDA流API中有多種溫和的方式來同步主機代碼。函數(shù)cudaStreamSynchronize(流)可以用于阻塞主機線程直到之前發(fā)布在指定流的所有操作完成為止悲伶。函數(shù)cudaStreamQuery(流)可以用于測試之前發(fā)布在指定流的所有操作是否完成艾恼，但不會阻塞主機線程。函數(shù)cudaEventSynchronize(事件)和cudaEventQuery(事件)與前兩種函數(shù)很像麸锉，區(qū)別在于后者是基于指定事件是否被記錄而前者是基于指定的流是否空閑钠绍。你也可以在一個單獨的流中基于一個特定的事件使用cudaStreamWaitEvent(事件)函數(shù)(即使事件被記錄在不同的流中或者不同的設備中！)

核函數(shù)執(zhí)行和數(shù)據(jù)傳輸?shù)闹丿B

之前我們已經(jīng)演示了如何在默認流中用主機端代碼來隱藏核函數(shù)執(zhí)行花沉。但是我們的主要目的是演示如何用核函數(shù)執(zhí)行隱藏數(shù)據(jù)傳輸柳爽。要實現(xiàn)它有幾點要求：

• 設備必須可以“并發(fā)地拷貝和執(zhí)行”。我們可以通過訪問cudaDeviceProp結構體的deviceOverlap屬性或者從CUDA SDK/Toolkit中deviceQuery示例程序的輸出中獲得碱屁。幾乎所有計算能力1.1及以上的設備都支持設備重疊磷脯。

• 核函數(shù)執(zhí)行和數(shù)據(jù)傳輸必須在不同的非默認流中。

• 涉及到數(shù)據(jù)傳輸?shù)闹鳈C內(nèi)存必須是固定主機內(nèi)存娩脾。

下面讓我們來修改上面的代碼以使用多個CUDA流赵誓，看一看是否實現(xiàn)了數(shù)據(jù)傳輸?shù)碾[藏。完整的代碼可以在Github上找到柿赊。在這個被修改的代碼中俩功，我們將大小為N的數(shù)組分為streamSize大小的數(shù)據(jù)塊。既然核函數(shù)可以獨立地操作所有數(shù)據(jù)碰声，那么每個數(shù)據(jù)塊也可以被獨立地處理诡蜓。流(非默認流)的數(shù)量nStreams=N/streamSize。實現(xiàn)數(shù)據(jù)的分解處理有多種方式胰挑，一種是將對每個數(shù)據(jù)塊的所有操作都放到一個循環(huán)中蔓罚，代碼如下所示：

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]);
}

另一種方式是將類似的操作放在一起批處理，首先發(fā)布所有主機到設備的數(shù)據(jù)傳輸洽腺，之后是核函數(shù)執(zhí)行脚粟，然后就是設備到主機的數(shù)據(jù)傳輸，代碼如下所示：

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]);
}

上述兩種異步方法都會產(chǎn)生正確的結果蘸朋，而且同一個流中相互依賴的操作都會按照需要的順序執(zhí)行。然而扣唱，這兩種方式的性能在不同版本的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

這里數(shù)據(jù)傳輸和核函數(shù)順序執(zhí)行的同步版本可以作為比較上述兩種異步版本是否有加速效果的基準炼彪。為什么這兩種異步執(zhí)行策略在不同架構上的效果不同呢吐根？為了解釋這一結果，我們需要了解CUDA設備如何調(diào)度和執(zhí)行任務辐马。CUDA設備中存在多種不同任務的引擎拷橘，它們會對發(fā)布的操作進行排隊。它們的功能就是維護不同引擎中任務間的依賴喜爷，但是在引擎內(nèi)部所有的外部依賴都會丟失冗疮；每個引擎中的任務都會按照它們被發(fā)布的順序執(zhí)行。C1060有一個單獨的拷貝引擎和一個單獨的核函數(shù)引擎檩帐。下圖是C1060運行上面示例代碼的時間線：

C1060 timelines

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

在這個原理圖中，我們假設主機到設備的數(shù)據(jù)傳輸湃密、核函數(shù)執(zhí)行诅挑、設備到主機三者所用的時間相同(所選擇的核函數(shù)代碼就是專門這樣設計的)。正如預料的那樣泛源，順序執(zhí)行的核函數(shù)并沒有任何操作重疊拔妥。對于異步版本1的代碼，拷貝引擎中的執(zhí)行順序是： H2D 1號流, D2H 1號流, H2D 2號流, D2H 2號流, 以此類推达箍。這就是為什么異步版本1沒有任何加速的原因：在拷貝引擎上任務的發(fā)布順序使得核函數(shù)執(zhí)行和數(shù)據(jù)傳輸無法重疊毒嫡。然而，從版本2較少的執(zhí)行時間來看幻梯，所有主機到設備的數(shù)據(jù)傳輸都在設備到主機的數(shù)據(jù)傳輸之前兜畸，是有可能實現(xiàn)重疊的。在原理圖中碘梢，我們可以看出異步版本理論時間是順序版本的8/12咬摇，前面的結果8.7ms剛好符合這個推算。

在C2050中煞躬，有兩個特征共同導致了它與C1060的性能差異肛鹏。C2050有兩個拷貝引擎，一個是用于主機到設備的數(shù)據(jù)傳輸恩沛，另一個用于設備到主機的數(shù)據(jù)傳輸在扰，第三個引擎是核函數(shù)引擎。下圖描述了C2050執(zhí)行示例代碼的時間線：

c2050 timelines

C2050具有兩個拷貝引擎恰好解釋了為什么異步版本1在C2050上具有很好的加速效果：與C1060正相反雷客，在stream[i]上設備到主機的數(shù)據(jù)傳輸并不會妨礙stream[i+1]上的主機到設備的數(shù)據(jù)傳輸芒珠，因為在C2050上每個方向的拷貝都有單獨的引擎。上面的原理圖顯示搅裙，該異步版本1的執(zhí)行時間大約是順序版本的一半皱卓，和實際結果相差無幾裹芝。

但是我們該如何解釋異步版本2在C2050上的性能下降呢？其實這與C2050可以并發(fā)執(zhí)行多個核函數(shù)有關娜汁。當多個核函數(shù)背靠背地被發(fā)布在不同的流(非默認流)中時嫂易，調(diào)度器會盡力確保這些核函數(shù)并發(fā)執(zhí)行，結果就導致每個核函數(shù)完成的信號被延遲掐禁，即所有核函數(shù)執(zhí)行完畢才發(fā)出信號怜械，而這個信號負責啟動設備到主機的數(shù)據(jù)傳輸。因此傅事，在異步版本2中缕允，主機到設備的數(shù)據(jù)傳輸與核函數(shù)執(zhí)行可以重疊，而核函數(shù)執(zhí)行與設備到主機的數(shù)據(jù)傳輸不能重疊享完。上面的原理圖中顯示異步版本2的總體時間大約是順序版本的9/12灼芭，正好與實驗結果7.5ms相吻合。

關于這個例子般又，在這篇文章CUDA Fortran Asynchronous Data Transfers中有更詳細的講解彼绷。讓人高興的是，對于計算能力3.5的設備(K20系列)茴迁，它所具有的超Q特性使得我們已經(jīng)不在需要特別安排啟動順序寄悯，所以上述兩個版本都會有很好的加速效果。我們會在將來的博客中討論如何使用開普勒的這些特性堕义。但是現(xiàn)在讓我們來看一下Tesla K20c GPU的運行結果猜旬。正如你所看到的，兩個異步執(zhí)行版本相比同步版本都有相同的加速效果倦卖。

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

總結

這篇文章和之前的文章都對如何優(yōu)化主機和設備間的數(shù)據(jù)傳輸進行了討論洒擦。之前的文章強調(diào)如何盡可能減少數(shù)據(jù)傳輸?shù)热蝿盏膱?zhí)行時間，這篇文章介紹了流以及如何使用它們來隱藏數(shù)據(jù)傳輸怕膛，即并發(fā)地執(zhí)行數(shù)據(jù)拷貝和核函數(shù)熟嫩。

說到流，我必須要提醒一點：盡管使用默認流非常的方便而且代碼寫起來也很簡單褐捻，但我們還是應該使用非默認流或者CUDA 7支持的每個線程單獨的默認流掸茅。尤其是在寫庫函數(shù)時，這一點尤為重要柠逞。如果在庫函數(shù)中使用默認流昧狮，那么對于庫函數(shù)用戶就不會有機會實現(xiàn)數(shù)據(jù)傳輸和核函數(shù)執(zhí)行的重疊了。

現(xiàn)在你應該明白了如何高效地在主機和設備間傳輸數(shù)據(jù)板壮，在下一篇博客中我們開始學習如何在核函數(shù)中高效的訪問數(shù)據(jù)逗鸣。