CUDA CC++ 最佳化資料傳輸效率

裝置記憶體和GPU之間的峰值頻寬(例如,在NVIDIA Tesla C2050上為144 GB / s)比主機記憶體和裝置記憶體之間的峰值頻寬(在PCIe x16 Gen2上為8 GB / s)高得多。這種差異意味著在主機和GPU裝置之間進行資料傳輸可能會影響或破壞整體的應用程式效能。

首先,讓我們瞭解主機-裝置之間資料傳輸的一些基本準則。

(1)儘可能減少在主機和裝置之間傳輸的資料量,即使這意味著與在主機CPU上執行相比,在GPU上執行核心的速度幾乎沒有提高。

(2)使用頁面鎖定(或“固定”)記憶體時,主機與裝置之間可能會有更高的頻寬。

(3)將許多小型傳輸分批成一個較大的傳輸的效能要好得多,因為它消除了大多數每次傳輸的開銷。

(4)有時,主機和裝置之間的資料傳輸可能會與核心執行和其他資料傳輸重疊。

1、使用nvprof測量資料傳輸時間

為了測量每次資料傳輸所花費的時間,我們可以在每次傳輸之前和之後記錄一個CUDA事件,並使用cudaEventElapsedTime()函式獲取所花費的時間。

但是,透過使用nvprof(CUDA工具包附帶的命令列CUDA探查器)(從CUDA 5開始),我們可以獲得經過的傳輸時間,而無需使用CUDA事件來對原始碼進行檢測。

具體示例如下:

CUDA C/C++ 最佳化資料傳輸效率

要分析此程式碼,我們使用nvcc對其進行編譯,然後使用程式檔名作為引數執行nvprof。

定位到cuda原檔案目錄下,使用

nvcc -o test。exe test。cu

來變異cuda程式。

使用

nvprof test。exe

執行cuda程式

當我在裝有GeForce GTX 1060的筆記本上執行時,得到以下輸出。

我所使用的平臺是Windows

CUDA C/C++ 最佳化資料傳輸效率

​如你所見,nvprof 測量每個CUDA memcpy呼叫所花費的時間。 它報告每個呼叫的平均,最小和最大時間。

2、最小化資料傳輸

我們不應該僅使用核心的GPU執行時間(相對於其CPU實現的執行時間)來決定執行GPU還是CPU版本。

我們還需要考慮在PCI-e總線上移動資料的成本,特別是在我們最初將程式碼移植到CUDA時。

由於CUDA的異構程式設計模型同時使用CPU和GPU,因此可以一次將程式碼移植到CUDA一個核心。

在移植的初始階段,資料傳輸可能會主導整個執行時間。將資料傳輸所花費的時間與核心執行所花費的時間分開進行記錄是值得的。

正如我們已經演示的,使用命令列事件探查器很容易。 隨著我們移植更多程式碼,我們將刪除中間傳輸並相應地減少總體執行時間。

3、固定主機記憶體

預設情況下,主機(CPU)資料分配是可分頁的。 GPU無法直接從可分頁的主機記憶體訪問資料,因此,當呼叫從可分頁的主機記憶體到裝置記憶體的資料傳輸時,CUDA程式必須先分配一個臨時的頁面鎖定或“固定”的主機陣列,才能複製主機資料到固定陣列,然後將資料從固定陣列傳輸到裝置記憶體,如下所示。

CUDA C/C++ 最佳化資料傳輸效率

​如圖所示,固定記憶體用作從裝置到主機的傳輸的暫存區域。

透過直接在固定記憶體中分配主機陣列,我們可以避免在可分頁和固定主機陣列之間進行傳輸的開銷。

使用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() 函式。

我們用以下的示例來比較可分頁和固定的傳輸速率。

CUDA C/C++ 最佳化資料傳輸效率

CUDA C/C++ 最佳化資料傳輸效率

​資料傳輸速率取決於主機系統的型別(主機板,CPU和晶片組)以及GPU。

在我的筆記本上執行,資料傳輸的頻寬基本上提升了2。5倍以上。

雖然這種方法可以提升速度,但是也不能過度地分配固定記憶體。 這樣做會降低整體系統效能,因為它會減少可用於作業系統和其他程式的物理記憶體量。

4、批處理小型資料傳輸

由於每次傳輸都會產生開銷,因此最好將許多小傳輸分批處理成單個傳輸。 這可以透過使用臨時陣列(最好是固定的)並將其與要傳輸的資料打包在一起來輕鬆實現。

對於二維陣列傳輸,可以使用cudaMemcpy2D()。

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

第一個引數:指目標資料的指標

第二個引數:目標陣列的間距

第三個引數:指向源資料的指標

第四個引數:源陣列的間距

第五個引數:要傳遞的子矩陣的寬度

第六個引數:要傳遞的子矩陣的高度

第七個引數:複製的方向(從主機到裝置或者從裝置到主機)

還有一個 cudaMemcpy3D() 函式用於傳輸3維陣列。

5、總結

主機和裝置之間的傳輸是GPU計算中涉及的資料移動的最慢環節,因此應注意儘量減少資料傳輸。遵循這篇文章中的準則可以幫助您確保必要的傳輸是有效的。建議不要使用帶有CUDA事件或其他計時器的工具程式碼來測量每次傳輸所花費的時間,建議使用nvprof,命令列CUDA分析器或一種視覺分析工具,例如NVIDIA Visual Profiler(也包括在內) 使用CUDA工具包。