1.
乙個gpu
上有很多的sm(
stream multiprocessor
),每個
sm中包括了8個
sp(stream processor
)標量流處理器,商業宣傳中所說的數百個「核」,大多指的是
sp的數量。隸屬於同乙個sm的
sp共用同一套取指與發射單元。
cuda
中的kernel
是以block
為單位執行的,乙個
block
必須在乙個
sm上執行,乙個
sp執行乙個執行緒,但是乙個
sm可以同時存在多個
block
的上下文。乙個
sm中活動執行緒塊的數量不超過
8個;所有活動執行緒塊中的
warp
數之和在計算能力
1.3裝置中,不超過32。
2.目前乙個
kernel
中只有乙個
grid
,grid
只能設定x、
y維,z維預設為
1,但是
block
的三維都可以自己設定。
3.乙個
block
中最多有
512個
thread
,這些thread
會被分割為執行緒束
warp在sm
上進行排程。每個
warp
包含連續的
32個執行緒,訪問儲存器時是以半執行緒束為單位的。
4.cuda
的原始檔必須使用
nvcc
編譯器進行編譯。這個編譯器在編譯的時候,會分離出主機端和裝置端**,主機端呼叫其他高效能編譯器編譯,如
gcc,裝置端**由
nvcc
編譯成ptx
**或者二進位制**。
ptx**是為動態編譯器
jit設計的,這樣可以應付不同顯示卡上的不同的機器語言。
5.在計算能力
1.3的裝置中,每個
sm的暫存器檔案數量為
16384
,每個暫存器檔案大小為
32bit
,如果每個執行緒使用的私有變數太多或者大小不定,可能將這樣的變數分配到區域性暫存器中(
local memory
)。例如
usigned int mt[3]
就是區域性暫存器的。
變數被分在**可以通過編譯輸出的
ptx彙編**檢視,有
.local
的則為區域性暫存器中的。
6.共享儲存器:
extren __shared__
宣告的變數都開始於同乙個位址,因此變數的布局必須通過顯式的定義,如
extern __shared__ char array;
__device__ void func()
7. cuda
主機端的記憶體被分為兩種:可分頁記憶體(
pageable memory
)(由malloc
和new
分配)和頁鎖定(
pinned
)記憶體(一定是儲存在物理記憶體,不會儲存在虛擬記憶體,因此速度較快,可以通過
dma加速與裝置通訊,使用
cudahostalloc
和cudafreehost
分配和釋放 )
cuda2.3
版本以上支援
portable memory
,在cudahostalloc
()時加上
cudahostallocportable
標誌,可以使多個
cpu共享一塊頁鎖定記憶體以實現
cpu間的通訊。預設情況下,
pinned
記憶體只能由分配的執行緒訪問。
write-combined memory
,在分配
cudahostalloc()
時設定cudahostallocwritecombined
的標誌,可以免除
cpu對記憶體的監視,減少將記憶體上的資料快取到l1、
l2 cache
中,在主機和視訊記憶體資料傳輸時節省了時間,這樣的作法比較適合於主機端只寫的資料。
主機端指標通過
cudahostalloc()
獲得,裝置端通過
cudahostgetdevicepointer()
獲得,這樣可以從
kernel
裡直接訪問主存,省略了在視訊記憶體上分配空間,實現了對記憶體的零拷貝訪問,不過使用前要通過
cudagetdeviceproperties()
判斷裝置是否有
canmaphostmemory
的屬性。
8.常數儲存:視訊記憶體上的唯讀位址空間,擁有快取加速,每個
sm擁有
8kb的常數儲存器快取。
__constant__ char p_hello;
可以通過
cudamemcpytosymbol(p_hello, hello_host, sizeof(char)*11)
從記憶體把資料拷進去。
9.紋理儲存器:唯讀儲存器,具備一些特殊的功能
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...