8. CUDA 記憶體使用 global 二------GPU的革命
序言:最近在另一個不寫技術的blog上,寫了最近的一些事情,或許是釋懷以後才會把心理面的事情寫出來,很感謝很多朋友能理解我現在的心情,有的朋友也會感到很驚訝,平時總看到我的時候都是很開心的樣子,很少會看到我不開心的時候,但是誰又會沒有煩惱的時候啦,……想想從大學到現在,經歷了很多,也是一直在思考,一直在反思,最開始的時候,想逃避,不過逃避沒用,學會了一個個問題的面對,一個一個的去淡然的接受,解決,釋懷。包容、釋懷、淡定、坦然,或許經歷再多的事情,就會更加的從容,更冷靜,有的人看著是成長的煩惱,或許我們更應該理解為成功的磨鍊,生活中有愛情,還有親情,友情,還有更多更多值得我們去體會,值得去思考的事情,去體會,去享受。當還有理想的時候,堅持……看近現代小小說的時候,最好旁邊放著佛經或者道德經,激情和包容不衝突,包容,有容乃大,或許更多的時候,冷靜的去思考,就像學習PMP的時候,專案經理做的事情,更多的時候是要聽,而不是說或者評論,釋懷,包容,在成長的路上,我也只是還在學習,還在體會,用心去感受,用心去思考。前言的內容或許有些太感性,但是真心喜歡更多的朋友能用心去體會,用心去思考問題的時候,問題其實不難。
正文:前面一章節已經寫到了記憶體訪問的問題,記憶體對齊的問題,不過在看到程式設計手冊第五章的時候,還是會有很多朋友問到我關於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,這個函式,注意使用~
文章來源:http://blog.csdn.net/OpenHero/archive/2008/12/15/3520578.aspx
來自 “ ITPUB部落格 ” ,連結:http://blog.itpub.net/22785983/viewspace-619845/,如需轉載,請註明出處,否則將追究法律責任。
相關文章
- NVDIA CUDA ---------GPU計算的革命GPU
- CUDA記憶體介紹記憶體
- CUDA一維紋理記憶體記憶體
- CUDA面記憶體用法總結記憶體
- 【CUDA學習】GPU硬體結構GPU
- CUDA 學習筆記之gpu結構筆記GPU
- 有Cuda能力的GPU核心GPU
- kaldi上使用gpu以及如何安裝cudaGPU
- GPU高效能程式設計CUDA實戰(二)GPU程式設計
- Tensorflow2對GPU記憶體的分配策略GPU記憶體
- 詳解GPU的記憶體頻寬與CPU的不同GPU記憶體
- Win10系統GPU共享記憶體怎麼關閉?Win10系統GPU共享記憶體的關閉方法Win10GPU記憶體
- CUDA 有 unified memory 還需要記憶體優化嗎?Nifi記憶體優化
- CUDA __global__ function 引數分析Function
- A6000單機多卡大模型訓練踩坑記錄(CUDA環境、多GPU卡住且視訊記憶體100%)大模型GPU記憶體
- aix的記憶體使用AI記憶體
- C語言-記憶體管理之二[記憶體指令]C語言記憶體
- PostgreSQLGPU加速(HeteroDBpg_strom)(GPU計算,GPU-DIO-NvmeSSD,列存,GPU記憶體快取)SQLGPU記憶體快取
- GPU體系架構(二):GPU儲存體系GPU架構
- JavaScript對記憶體的使用JavaScript記憶體
- RMAN 對記憶體的使用記憶體
- Linux共享記憶體(二)Linux記憶體
- Windows記憶體管理分析(二)Windows記憶體
- oracle 記憶體結構(二)Oracle記憶體
- MySQL記憶體使用MySql記憶體
- ubuntu解決GPU視訊記憶體佔用問題UbuntuGPU記憶體
- LightGCN實踐2——GPU記憶體爆炸終結篇GCGPU記憶體
- [Virtualization]ESXi體系結構與記憶體管理(二)控制記憶體分配記憶體
- Oracle記憶體分配與使用小記(二)Shared Pool and Large PoolOracle記憶體
- GPU深度學習效能的三駕馬車:Tensor Core、記憶體頻寬與記憶體層次結構GPU深度學習記憶體
- nginx中共享記憶體的使用Nginx記憶體
- Oracle的記憶體分配和使用Oracle記憶體
- linux記憶體管理(二)- vmallocLinux記憶體
- CUDA(五)用deviceQuery看GPU屬性devGPU
- 開源大模型佔GPU視訊記憶體計算方法大模型GPU記憶體
- Linux記憶體使用的體會(轉)Linux記憶體
- GPU的並行運算與CUDA的簡介GPU並行
- Java的記憶體 -JVM 記憶體管理Java記憶體JVM