本文翻譯自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
正如咱們在前面文章中所說的那樣,咱們能夠在數據傳輸的先後使用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
來編譯,而後以程序的名字爲參數運行nvprof
。app
$ 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 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事件或者其餘計時器函數。
這篇文章的重點是如何高效地進行數據傳輸。下一篇文章,咱們會討論如何用計算和其餘數據傳輸來隱藏數據傳輸。