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