參考:
硬體上,gpu由多個sm(steaming multiprocessor)構成,sm有多個warp,warp有多個sp(streaming processor),乙個sp對應乙個執行緒。
乙個warp中的sp執行相同的指令。
block內部可以使用sm提供的shared memory和__syncthreads()功能實現執行緒同步和通訊。但是block之間除了結束kernel之外是無法同步的。
由於乙個warp中的sp執行相同的指令,所以如果**中有分支,而乙個warp中一部分執行緒執行乙個分支,一部分執行緒執行另乙個分支,則會造成一些執行緒停滯。所以要盡量保證乙個warp內的執行緒都執行同乙個分支。
由於block是放到乙個sm中執行的,而乙個sm中的所有sp都共享乙個l1 cache,所以如果乙個block中的執行緒操作的記憶體位址相距很遠,那麼一些cache line被load到l1 cache後沒有怎麼被用到就被其他執行緒load進來的cache line擠出去了,那就會造成嚴重的讀寫放大,嚴重影響效率。
GPU程式設計筆記(2)
2009 02 02 19 52 高階渲染語言基礎 1 hlsl語法與c語法非常類似。2 資料型別 bool int 32位signed half 16位float float 32位float double 64位float 3 變數宣告 與c一樣 4 型別修飾 可以使用const,與c 一樣 r...
GPU 程式設計入門到精通(四)之 GPU 程式優化
gpu 程式設計入門到精通 三 之 第乙個 gpu 程式 中講到了怎樣利用 cuda5.5 在 gpu 中執行乙個程式。通過程式的執行。我們看到了 gpu 確實能夠作為乙個運算器。可是,我們在前面的樣例中並沒有正真的發揮 gpu 並行處理程式的能力。也就是說之前的樣例僅僅利用了 gpu 的乙個執行緒...
GPU程式設計 四 並行規約優化
如果之前沒有用過gdb,可以速學一下,就幾個指令.想要用cuda gdb對程式進行除錯,首先你要確保你的gpu沒有在執行作業系統介面,比方說,我用的是ubuntu,我就需要用sudo service lightdm stop關閉圖形介面,進入tty1這種字元介面.當然用ssh遠端訪問也是可以的.接下...