cuda學習筆記(一)儲存

2021-06-18 04:00:06 字數 2613 閱讀 6534

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...