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() { short* array0 = (short*)array; float* array1 = (float*)&array0[128]; int array2 = (int*)&array1[64]; } 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中，在主機和視訊記憶體資料傳輸時節省了時間，這樣的作法比較適合於主機端只寫的資料。 mapped memory,主機端指標通過cudaHostAlloc()獲得，裝置端通過cudaHostGetDevicePointer()獲得，這樣可以從kernel裡直接訪問主存，省略了在視訊記憶體上分配空間，實現了對記憶體的零拷貝訪問，不過使用前要通過cudaGetDeviceProperties()判斷裝置是否有canMapHostMemory的屬性。 8.常數儲存：視訊記憶體上的只讀地址空間，擁有快取加速，每個sm擁有8KB的常數儲存器快取。 __constant__ char p_hello[];  可以通過cudaMemcpyToSymbol(p_hello, hello_host, sizeof(char)*11)從記憶體把資料拷進去。 9.紋理儲存器：只讀儲存器，具備一些特殊的功能。