OpenCL memory object 之 傳輸優化

首先我們了解一些優化時候的術語及其定義：

?

1、deferred allocation（延遲分配），

???? 在第一次使用memory object傳輸數據時，runtime才對memory object真正分配空間。 這樣減少了資源浪費，但第一次使用時要慢一些[一個context多個設備，一個memory object多個location，見前面的blog]。

?

?

2.peak interconntect bandwith(峰值內聯帶寬）

???? host和device之間通過PCIE總線傳輸數據，PCIE2.0的上行、下行帶寬都是8Gb/s, 對于我們的程序，能達到3Gb/s就不錯了，我的筆記本測試只有1.2Gb/s。

?

?

3.Pinning(對內存實施pinning操作）

???? host memory準備向gpu傳輸時，都要首先進行pinning，就是lock page（禁止交換到外存），pinning操作有一定的性能開銷，開銷的大小和pinning的host memory大小有關，越大就開銷越大。我們可以把host memory分配到pre pinned memory中減少這種開銷。

?

4.WC(write combined operation)

?? WC是cpu寫固定地址時的一個特性，通過把鄰接的寫操作綁定到一個cacheline，然后發一個寫請求，實現了批量寫操作。[Gpu內部也有相似的地址合并操作]

?

?

5.uncached access

??? 一些內存區域被配置為uncache access，cpu訪問比較慢，但是有利于向device memory傳輸數據，比如前篇日志提到device visible host memory。

?

?

6.USWC(無cache的寫綁定）

?? gpu訪問uncached的host memory不會產生cache一致性問題，速度會比較快，cpu寫因為WC也比較快，相對來說cpu讀會變慢。在APU上，這個操作會提供一個快速的cpu寫，gpu讀的path。

?

下面看看buffer的分配及使用:

?

1.normal buffer

????? 用CL_MEM_READ_ONLY/CL_MEM_WRITE_ONLY/CL_MEM_READ_WRITE標志創建的buffer位于device memory中，GPU能夠以很高的bandwidth訪問這些它，例如對一些高端的顯卡，超過100GB/s,host要訪問這些內存，只能通過peak interconntect bandwith（PCIE)。

?

2. zero copy buffer

???? 這種buffer并不做實際的copy工作（除非特殊指定執行copy操作，比如clEnqueueCopyBuffer）。根據創建buffer的type參數，它可能位于host memory也可能位于device memory。

?

???? 如果device及操作系統支持zero copy，則下面buffer類型可以使用:

?

? The CL_MEM_ALLOC_HOST_PTR buffer
– zero copy buffer駐留在host。
– host能夠以全帶寬訪問它。
– device通過interconnect bandwidth訪問它。
– 這塊buffer被分配在prepinned的host memory中。

? The CL_MEM_USE_PERSISTENT_MEM_AMD buffer is
– zero copy buffer 駐留在GPU device中。
– GPU能全帶寬訪問它。
– host能夠以interconnect帶寬訪問它 (例如streamed寫帶寬host->device，低的讀帶寬，因為沒有cache利用）。

– 在host和device之間通過interconnect帶寬傳輸數據。

?

注意：創建buffer的大小是平臺dependience的，比如在某個平臺上一個buffer不能超過64M，總的buffer不能超過128M等。

?

zero copy內存在APU上可以得到很好的效果，cpu可以高速的寫，gpu能夠高速的讀，但因為無cache，cpu讀會比較慢。

1. buffer = clCreateBuffer(CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY)
2. address = clMapBuffer( buffer )
3. memset( address ) or memcpy( address ) (if possible, using multiple CPU
cores)
4. clEnqueueUnmapMemObject( buffer )
5. clEnqueueNDRangeKernel( buffer? )

對于數據量小的傳輸，zero copy時延（map，unmap等）通常低于相應的DMA引擎時延。

?

3. prepinned buffer

???? pinned buffer類型是CL_MEM_ALLOC_HOST_PTR/CL_MEM_USE_HOST_PTR, buffer初始就被創建在prepinned內存中。 EnqueueCopyBuffer以interconnect帶寬在host和device之間傳輸數據（沒有pinned和unpinned開銷）。

??? 注意：CL_MEM_USE_HOST_PTR能夠把已經存在的host buffer轉化到pinned memory中去，但是為了保證傳輸速度，host buffer必須保證256字節對齊。如果只是用來傳輸數據的話，CL_MEM_USE_HOST_PTR 類型memory對象會一直為prepinned內存，但是它不能作為kernel參數。如果buffer要在kernel中使用的話，runtime會在device創建一個該buffer cache copy，接下來的copy操作不會通過fast path（要保持cache一致性）。

?

下面的一些函數支持prepinned memory，注意：讀取memory可以使用offset：
? clEnqueueRead/WriteBuffer
? clEnqueueRead/WriteImage
? clEnqueueRead/WriteBufferRect (Windows only)