OpenCL memory object 之 传输优化
首先我們了解一些優(yōu)化時候的術語及其定義：
?
1、deferred allocation（延遲分配），
???? 在第一次使用memory object傳輸數(shù)據(jù)時，runtime才對memory object真正分配空間。 這樣減少了資源浪費，但第一次使用時要慢一些[一個context多個設備，一個memory object多個location，見前面的blog]。
?
?
2.peak interconntect bandwith(峰值內(nèi)聯(lián)帶寬）
???? host和device之間通過PCIE總線傳輸數(shù)據(jù)，PCIE2.0的上行、下行帶寬都是8Gb/s, 對于我們的程序，能達到3Gb/s就不錯了，我的筆記本測試只有1.2Gb/s。
?
?
3.Pinning(對內(nèi)存實施pinning操作）
???? host memory準備向gpu傳輸時，都要首先進行pinning，就是lock page（禁止交換到外存），pinning操作有一定的性能開銷，開銷的大小和pinning的host memory大小有關，越大就開銷越大。我們可以把host memory分配到pre pinned memory中減少這種開銷。
?
4.WC(write combined operation)
?? WC是cpu寫固定地址時的一個特性，通過把鄰接的寫操作綁定到一個cacheline，然后發(fā)一個寫請求，實現(xiàn)了批量寫操作。[Gpu內(nèi)部也有相似的地址合并操作]
?
?
5.uncached access
??? 一些內(nèi)存區(qū)域被配置為uncache access，cpu訪問比較慢，但是有利于向device memory傳輸數(shù)據(jù)，比如前篇日志提到device visible host memory。
?
?
6.USWC(無cache的寫綁定）
?? gpu訪問uncached的host memory不會產(chǎn)生cache一致性問題，速度會比較快，cpu寫因為WC也比較快，相對來說cpu讀會變慢。在APU上，這個操作會提供一個快速的cpu寫，gpu讀的path。
?
下面看看buffer的分配及使用:
?
1.normal buffer
????? 用CL_MEM_READ_ONLY/CL_MEM_WRITE_ONLY/CL_MEM_READ_WRITE標志創(chuàng)建的buffer位于device memory中，GPU能夠以很高的bandwidth訪問這些它，例如對一些高端的顯卡，超過100GB/s,host要訪問這些內(nèi)存，只能通過peak interconntect bandwith（PCIE)。
?
2. zero copy buffer
???? 這種buffer并不做實際的copy工作（除非特殊指定執(zhí)行copy操作，比如clEnqueueCopyBuffer）。根據(jù)創(chuàng)建buffer的type參數(shù)，它可能位于host memory也可能位于device memory。
?
???? 如果device及操作系統(tǒng)支持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帶寬傳輸數(shù)據(jù)。
?
注意：創(chuàng)建buffer的大小是平臺dependience的，比如在某個平臺上一個buffer不能超過64M，總的buffer不能超過128M等。
?
zero copy內(nèi)存在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? )
對于數(shù)據(jù)量小的傳輸，zero copy時延（map，unmap等）通常低于相應的DMA引擎時延。
?
3. prepinned buffer
???? pinned buffer類型是CL_MEM_ALLOC_HOST_PTR/CL_MEM_USE_HOST_PTR, buffer初始就被創(chuàng)建在prepinned內(nèi)存中。 EnqueueCopyBuffer以interconnect帶寬在host和device之間傳輸數(shù)據(jù)（沒有pinned和unpinned開銷）。
??? 注意：CL_MEM_USE_HOST_PTR能夠把已經(jīng)存在的host buffer轉化到pinned memory中去，但是為了保證傳輸速度，host buffer必須保證256字節(jié)對齊。如果只是用來傳輸數(shù)據(jù)的話，CL_MEM_USE_HOST_PTR 類型memory對象會一直為prepinned內(nèi)存，但是它不能作為kernel參數(shù)。如果buffer要在kernel中使用的話，runtime會在device創(chuàng)建一個該buffer cache copy，接下來的copy操作不會通過fast path（要保持cache一致性）。
?
下面的一些函數(shù)支持prepinned memory，注意：讀取memory可以使用offset：
? clEnqueueRead/WriteBuffer
? clEnqueueRead/WriteImage
? clEnqueueRead/WriteBufferRect (Windows only)
關注 - 7
粉絲 - 127 +加關注 1 0 (請您對文章做出評價) ? 上一篇：OpenCL memory object 之 Global memory (2)
? 下一篇：OpenCL memory object 之選擇傳輸path
posted on 2011-12-18 13:52 邁克老狼2012 閱讀(841) 評論(1) 編輯 收藏
評論
#1樓??
老狼，最近一直在你的博客上學習，學到了很多東西，這篇文章有點東西不太明白，幫忙解答下唄1、在某個平臺上一個buffer不能超過64M，總的buffer不能超過128M等
總的buffer是指CPU+GPU的zero-copy buffer嗎 ， 分別64M？
2、clMapBuffer( buffer ) 是必須的嗎 它和clEnqueueMapBuffer有什么區(qū)別？用clEnqueueMapBuffer可以代替clmMapBuffer嗎
總結
以上是生活随笔為你收集整理的OpenCL memory object 之 传输优化的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： IOS开发之MD5加密和钥匙串的使用-o
- 下一篇： AFN框架和SDWebImage框架的上