CUDA的儲存器結構

2021-09-29 00:22:52 字數 992 閱讀 7006

因為寫部落格主要想記錄比較感興趣的點,所以可能順序不那麼通常、、通暢。

cuda c有共享記憶體(shared memory),它的關鍵字是__share__。新增了該關鍵字的變數表明該變數是乙個共享變數。

共享記憶體被乙個執行緒塊內的所有執行緒共享,塊內線程可以修改它的值,它的生命週期和它所在的執行緒塊一樣。

因為共享記憶體的緩衝區在物理gpu上,而不是gpu外的系統記憶體上,所以訪問共享記憶體的延遲要低於訪問普通緩衝區的延遲,也就是使用共享記憶體可以提速。

使用shared memory時,需要注意執行緒的同步問題。當所有執行緒都進行完了對共享記憶體的某一操作後(比如寫操作),才能執行後面的語句(比如從共享記憶體讀取資料),否則結果會出錯。所以需要將執行緒同步,讓快的執行緒等一下慢的。執行緒同步可以使用__syncthreads()語句。

先放一段《cudy by example》的向量求和來看一下共享記憶體的使用:

__global__ void

dot(

float

*a,float

*b,float

*c)__syncthreads()

;//防止還有執行緒沒算完,先等其他執行緒算完了再進行後面的操作..

..}

除了全域性記憶體,cuda定義了常量記憶體,用於儲存核函式執行期間不會發生變化的數值(比如輸入的資料)。

常量記憶體的關鍵字為__constant__,比如

__constant__ float s[30]

與全域性記憶體相比,常量記憶體分配時用malloc(),從主機拷貝記憶體到裝置時使用cudamemcpytosymbol()而不是cudamemcpy()。

常量記憶體的好處有啥?

常量記憶體的一次讀取操作可以擴散到半個執行緒束(16個執行緒),也就是假設這16個執行緒要求訪問常量記憶體的相同位置的資料時,只需一次操作。

常量記憶體被快取在gpu上,對相同位置的連續讀操作時,命中快取

CUDA學習筆記之 CUDA儲存器模型

gpu片內 register,shared memory 板載視訊記憶體 local memory,constant memory,texture memory,texture memory,global memory host 記憶體 host memory,pinned memory.cuda儲...

CUDA學習筆記之 CUDA儲存器模型

標籤 cuda 儲存binding cache 程式設計api 2010 01 03 20 32 5577人閱讀收藏 舉報 cuda 6 cuda儲存器模型 gpu片內 register,shared memory 板載視訊記憶體 local memory,constant memory,textu...

CUDA學習筆記之 CUDA儲存器模型

標籤 cuda 儲存binding cache 程式設計api 2010 12 14 01 33 1223人閱讀收藏 舉報 cuda 26 gpu片內 register,shared memory 板載視訊記憶體 local memory,constant memory,texture memory...