CUDA的global内存访问的问题
生活随笔
收集整理的這篇文章主要介紹了
CUDA的global内存访问的问题
小編覺得挺不錯的,現在分享給大家,幫大家做個參考.
http://blog.csdn.net/OpenHero/article/details/3520578
關于CUDA的global內存訪問的問題,怎么是訪問的沖突,怎樣才能更好的訪問內存,達到更高的速度。下面先看幾張圖,這些圖都是CUDA編程手冊上的圖,然后分別對這些圖做解釋,來理解硬件1.0,1.1 以及現在最新的硬件的訪問內存的區別。
我們在這里再深入的講解一下global內存對齊的問題,每次執行一條明命令的時候,都是會按照32個thread為一個warp,一起來執行,但是在執行的時候,又會按照硬件的條件(這里有兩個限制條件,一個是內存訪問的時鐘和執行core的時鐘不一樣,第二個是為了細粒度的分支的問題)然后就會把16個thread組成的half-warp來一次訪問global內存才能讓訪問內存的性能高一些,這個可以理解;
就像手冊上說的那樣,如果16個thread(half-warp)訪問內存的時候,如果每一個thread訪問32bits就是4個字節,那么就可以合并為一個64bytes的訪問,手冊上這點寫得有點讓人咋一看不太明白~4bytes(32bits)*16 = 64bytes,就是這么來的,如果每一個thread訪問64bits(8個bytes),那么就可以合并為128bytes的訪問;這里啦,由于合并訪問的最大限制是128bytes,所以最大也按照128bytes一次訪問來合并,如果超過,就得多次訪問,或者如果沒有按照這樣的方式對齊訪問,也會多次訪問;下面是1.2device之前的訪問的幾個圖,這里要把1.2device以前和以后的分開,是因為這里在對齊訪問的方式的時候,有不同的策略;先看1.2device以前的能合并為一次訪問的情況:
下圖是編程手冊上的圖:
這里的每一個thread都是訪問的對應的地址,是對齊的,所以可以合并為一個存儲event;
下面這個圖是沒有對齊訪問,就造成了non-coalesced訪問的問題,下面可以看圖說話:
左邊的那個好理解,中間對應的thread訪問的地址交叉了,thread3和thread4交叉訪問了,在硬件1.2版本之前的都會造成Non-Coalesced訪問;
詳細的需要說明的是右邊的為什么也造成了Non-Coalesced(非對齊)訪問,這個是基礎問題,大家理解的內存對齊是怎么樣的?按照固定思路,或者教材上強調的都是中間過程的內存訪問的對齊,但是內存是從offset 0x00000000位置開始的,就是偏移量0開始的,如果要真的滿足內存對齊,嚴格的說起來就需要從內存的0地址開始算起,再加上我們知道的global內存的對齊方式有幾種,4位bytes,8bytes,16bytes,這里說的對齊方式,注意區別關系;再來看看右邊的那個圖:thread0開始,從address128的位置開始向下便宜的位置是?132-128=4 偏移了4,16個threads整體訪問的是16*4=64,是從132開始的,從0算起來,132-0 =132; 132/16 = 8…4,從整體上講,從0偏移位置開始,偏移了4個位置,這里的就造成了訪問的未對齊,這個是從整體角度上講的,和左邊的圖比較一下,那個是按照局部對齊來說的,注意理解;
繼續看圖說話:
左邊的圖看看,算一下,局部的時候偏移了,從局部和整體來說,都會引起未對齊訪問;
右邊的圖自己算一下,是不是超出了剛才我說的范圍;所以造成了內存訪問的未對齊情況;
前面我們看的圖都是1.2版本前的硬件的情況下的內存訪問情況,現在看看1.2版本以后的硬件;
這里解釋一下,什么叫1.2版本的硬件,或許有些朋友也不太了解,g80架構的都是1.0或者1.1的硬件架構,現在的gtx200系列的都是1.3的架構,其實1.2的硬件架構,或許是Nvidia的一個內部的,沒有推出產品,可能準備提供給低端的產品,但是我想沒有推出低端的產品,直接就上1.3device了,市場需求吧~~如果下一步GTX的架構還是按照老路子,不改進的話,或許Intel的Lrb上來以后,對Nvidia的產品,就是一個很大的競爭了;
不說廢話了,先看圖:
1.2以后的硬件版本,弱化了threads之間交叉訪問的時候,沒對齊的情況,只要大家都在一次訪問的64bytes的一個段里面,或者128bytes的一個段里面面,這樣的段訪問,那就可以不用多次訪問,當然如果你16個threads分別跨過了16個段,那就得產生16個存儲event~記住幾個段的定義,這里說的段,就是我們常常理解的對齊的方式,全局的內存訪問對齊方式,8個bits的是按照32bytes對齊,16bits的是按照64bytes對齊,32bits和64bits都是按照128bytes對齊;
在優化代碼的時候,這個地方是一個值得注意的部分;
API函數里面有對齊訪問的接口,會按照對齊的方式分配global內存給你,不過注意其中的一個offset值的使用,這個是為了解決全局情況下的對齊偏移的問題:)cudaMallocPitch,這個函數,注意使用~
總結
以上是生活随笔為你收集整理的CUDA的global内存访问的问题的全部內容,希望文章能夠幫你解決所遇到的問題。
- 上一篇: extern __shared__
- 下一篇: __syncthreads()