CUDA程式設計——cudaHostAlloc

CUDA程式設計——zero copy

2016年02月04日 13:40:38 ZhangJunior 閱讀數：4593

 版權聲明：本文為部落客原創文章，未經部落客允許不得轉載。 https://blog.csdn.net/junparadox/article/details/50633641

零複制

　　zero copy（零複制）是一種特殊形式的記憶體映射，它允許你将host記憶體直接映射到裝置記憶體空間上。其實就是裝置可以通過直接記憶體通路（direct memory access，DMA）方式來通路主機的鎖頁記憶體。 

　　

鎖頁主機記憶體

　　現代作業系統都支援虛拟記憶體，作業系統實作虛拟記憶體的主要方法就是通過分頁機制。作業系統将記憶體中暫時不使用的内容換出到外存（硬碟等大容量存儲）上，進而騰出空間存放将要調入記憶體的資訊。這樣，系統好像為使用者提供了一個比實際記憶體大得多的存儲器，稱為虛拟存儲器。 

　　鎖頁就是将記憶體頁面标記為不可被作業系統換出的記憶體。是以裝置驅動程式給這些外設程式設計時，可以使用頁面的實體位址直接通路記憶體（DMA），進而避免從外存到記憶體的複制操作。CPU 仍然可以通路上述鎖頁記憶體，但是此記憶體是不能移動或換頁到磁盤上的。CUDA 中把鎖頁記憶體稱為pinned host memory 或者page-locked host memory。

鎖頁主機記憶體的優勢

　　使用鎖頁記憶體（page-locked host memory）有一些優勢：

• 鎖頁記憶體和GPU記憶體之間的拷貝可以和核心程式同時執行，也就是異步并發執行。
• 在一些裝置上鎖頁記憶體的位址可以從主機位址空間映射到CUDA 位址空間，免去了拷貝開銷。
• 在擁有前線總端的系統上，如果主機記憶體被配置設定為鎖頁記憶體，主機記憶體和GPU 記憶體帶寬可以達到更高，如果主機記憶體被配置設定為Write-Combining Memory，帶寬會進一步提升。

然而鎖頁主機存儲器是稀缺資源，是以鎖頁記憶體配置設定得多的話，配置設定會失敗。另外由于減少了系統可分頁的實體存儲器數量，配置設定太多的分頁鎖定記憶體會降低系統的整體性能

使用鎖頁主機記憶體

　　在GPU 上配置設定的記憶體預設都是鎖頁記憶體，這隻是因為GPU 不支援将記憶體交換到磁盤上。在主機上配置設定的記憶體預設都是可分頁，如果需要配置設定鎖頁記憶體，則需要使用cudaMallocHost() 或者cudaHostAlloc()。釋放時需要使用cudaFreeHost() 釋放這一塊記憶體。調用正常的C函數釋放，可能會崩潰或者出現一些不常見的錯誤。也可以通過函數cudaHostRegister() 把可分頁記憶體标記為鎖頁記憶體。

__host__ ​cudaError_t cudaMallocHost ( void** ptr, size_t size )

__host__ ​cudaError_t cudaHostAlloc ( void** pHost, size_t size, unsigned int  flags )

__host__ ​cudaError_t cudaFreeHost ( void* ptr )
           

cudaHostAlloc() 多了一個可選形參flags ，功能更強大。flags 的值可以取如下值。

#define cudaHostAllocDefault 0x00
Default page-locked allocation flag

#define cudaHostAllocMapped 0x02
Map allocation into device space

#define cudaHostAllocPortable 0x01
Pinned memory accessible by all CUDA contexts

#define cudaHostAllocWriteCombined 0x04
Write-combined memory
           

cudaHostRegister() 函數用于把已經的存在的可分頁記憶體注冊為分頁鎖定的。

__host__ ​cudaError_t cudaHostRegister ( void* ptr, size_t size, unsigned int  flags )
           

• 1

flags 是一個可選形參，可以取如下值。

#define cudaHostRegisterDefault 0x00
Default host memory registration flag

#define cudaHostRegisterIoMemory 0x04
Memory-mapped I/O space

#define cudaHostRegisterMapped 0x02
Map registered memory into device space

#define cudaHostRegisterPortable 0x01
Pinned memory accessible by all CUDA contexts
           

下面分别介紹這些flags 的作用。

Portable Memory

　　一塊鎖頁記憶體可被系統中的所有裝置使用（一個系統中有多個CUDA裝置時）。 啟用這個特性需要在調用cudaHostAlloc() 時使用cudaHostAllocPortable 選項，或者在調用cudaHostRegister() 使用cudaHostRegisterPortable 選項。 

　　

Write-Combining Memory

　　預設情況下，鎖頁主機存儲是可緩存的。可以在調用cudaHostAlloc() 時傳入cudaHostAllocWriteCombined 标簽使其被配置設定為寫結合的（Write-Combining Memory）。寫結合存儲不使用L1 和L2 cache，是以程式的其它部分就有更多的緩存可用。此外，寫結合記憶體通過PCI-E 傳輸資料時不會被監視（snoop），這能夠獲得高達40%的傳輸加速。 從主機讀取寫結合存儲非常慢（因為沒有使用L1、L2cache），是以寫結合存儲應當隻用于那些主機隻寫的存儲。 

　　

Mapped Memory

　　一塊鎖頁記憶體可以在調用cudaHostAlloc() 配置設定時傳入cudaHostAllocMapped 标簽或者在使用cudaHostRegister() 注冊時使用cudaHostRegisterMapped 标簽，把鎖頁記憶體位址映射到裝置位址空間。這樣，這塊存儲會有兩個位址：一個是從cudaHostAlloc() 或malloc() 傳回的在主機記憶體位址空間上；另一個在裝置存儲器上，可以通過cudaHostGetDevicePointer() 取得。核心函數可以使用這個指針通路這塊存儲。 cudaHostAlloc() 傳回的位址指針一個的例外情況是，主機和裝置使用統一位址空間（Unified Virtual Address Space）。 

核心直接存取主機記憶體有很多優勢：

• 無需在裝置上配置設定記憶體，也無需在主機記憶體和裝置記憶體之間拷貝資料。資料傳輸是在核心需要的時候隐式進行的。
• 無須使用流（cuda stream）就可以并發資料傳輸和核心執行；資料傳輸和核心執行自動并發執行。

因為映射的鎖頁主機記憶體是主機和裝置之間共享的，是以在使用cuda stream 或者cuda event 時必須對記憶體讀寫同步；避免潛在的寫後讀，讀後寫或者寫後寫等多線程同步問題。 

　　為了能夠對任何映射的鎖頁主機記憶體解引用裝置指針，必須在調用任何cuda 運作時函數前調用cudaSetDeviceFlags()，并傳入cudaDeviceMapHost 标簽。否則，cudaHostGetDevicePointer() 将會傳回錯誤。 

　　如果裝置不支援被映射分頁鎖定存儲，cudaHostGetDevicePointer() 将會傳回錯誤。程式員可以檢查canMapHostMemory 屬性，如果裝置支援映射鎖頁主機記憶體，将會傳回1。

注意：使用映射鎖頁主機記憶體看，原子操作将不再保證原子性。cudaHostRegisterIoMemory 是cudaHostRegister() 特有的選項，可以把主機記憶體映射到IO 位址空間。

參考文獻

[1]https://en.wikipedia.org/wiki/CUDA_Pinned_memory 

[2] Cook, Shane (2013). CUDA Programming: A Developer’s Guide to Parallel Computing with GPUs (1st ed.). Morgan Kaufmann Publishers Inc. pp. 334–335. ISBN 9780124159334.