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)