常量記憶體，顧名思義，它是唯讀記憶體。這種型別的記憶體要麼是在編譯時宣告為唯讀記憶體，要麼是在執行時通過主機端定義為唯讀記憶體。常量只是從gpu記憶體的角度而言。在編譯時宣告一塊常量記憶體，需要用到const關鍵字。
常量記憶體其實只是全域性記憶體的一種虛擬位址形式，並沒有特殊保留的常量記憶體塊。常量記憶體有兩個特性，乙個是快取記憶體，另乙個是它支援將單個值廣播到執行緒束中的每個執行緒。但要注意的是，對於那些資料不太集中或者資料重用率不高的記憶體訪問，盡量不要使用常量記憶體。
當常量記憶體將資料分配或廣播到執行緒束中的每個執行緒時（注意，實際上硬體會將單次記憶體讀取操作廣播到半個執行緒束），廣播能夠在單個週期內發生，因此這個特性是非常有用的。雖然當所有16個執行緒都讀取相同位址時，這個功能可以極大提高效能，但當所有16個執行緒分別讀取不同的位址時，它實際上會降低效能。如果半個執行緒束中的所有16個執行緒需要訪問常量記憶體中的不同資料，那麼這個16次不同的讀取操作會被序列化，從而需要16倍的時間來發出請求。但如果從全域性記憶體中讀取，那麼這些請求就會同時發出。這種情況下，從常量記憶體讀取就會慢於從全域性記憶體中讀取。
需要注意的是，當我們宣告乙個核心常量的時候，在編譯器將cuda c**轉換成ptx彙編**時會用字面值（0x55555555）直接替換常量值(data)的位址。
const int data = 0x55555555;
int d = data; //此時data會直接編譯為字面值0x55555555
const int data[3] = ;
int d = data[1]; //此時data[1]會被編譯為data[1]的位址
gpu上的常量記憶體並不是真正意義上的常量記憶體，因為gpu沒有專門常量記憶體預留的特殊記憶體區。由於常量記憶體時通過16位的位址進行訪問的，而16位位址能夠快速進行訪問，因此常量記憶體最大限制為64kb。在對常量記憶體進行資料更新時，需要在主機端呼叫cudamemcpytosymbol函式。
int cpu_data[512] = 0;
__constant__ const
int gpu_const_data[512];
__device__ static
int gpu_static_data[512];
cudamemcpytosymbol(gpu_const_data, cpu_data, 512 * sizeof(int));
cudamemcpytosymbol(gpu_static_data, cpu_data, 512 * sizeof(int));
CUDA 常量記憶體
常量記憶體的大小限定為64k，每個sm擁有8kb的常數儲存器快取，在編譯期時宣告一塊常量記憶體，需要用到 constant 關鍵字，例如 constant float my array 1024 不同於c c 中的const常量，cuda中常量內存在宣告後是可以修改的。如果要在執行時改變常量記憶體中...
CUDA學習之零拷貝記憶體
當使用零拷貝記憶體來共享主機和裝置間的資料時，必須同步主機和裝置間的記憶體訪問，同時更改主機和裝置的零拷貝記憶體中的資料將導致不可預知的後果。有兩種常見的異構計算系統架構 整合架構和離散架構。在整合架構中，cpu和gpu整合在乙個晶元上，並且在實體地址上共享主存。在這種架構中，由於無須在pcie匯流...
CUDA學習 GPU記憶體
通俗意義上的裝置記憶體 位置 裝置記憶體。形式 關鍵字 shared 新增到變數宣告中。如 shared float cache 10 目的 對於gpu上啟動的每個執行緒塊，cuda c編譯器都將建立該共享變數的乙個副本。執行緒塊中的每個執行緒都共享這塊記憶體，但執行緒卻無法看到也不能修改其他執行緒...