1.常量記憶體

常量記憶體，顧名思義，它是隻讀記憶體。這種型別的記憶體要麼是在編譯時宣告為只讀記憶體，要麼是在執行時通過主機端定義為只讀記憶體。常量只是從GPU記憶體的角度而言。在編譯時宣告一塊常量記憶體，需要用到const關鍵字。
常量記憶體其實只是全域性記憶體的一種虛擬地址形式，並沒有特殊保留的常量記憶體塊。常量記憶體有兩個特性，一個是快取記憶體，另一個是它支援將單個值廣播到執行緒束中的每個執行緒。但要注意的是，對於那些資料不太集中或者資料重用率不高的記憶體訪問，儘量不要使用常量記憶體。

當常量記憶體將資料分配或廣播到執行緒束中的每個執行緒時（注意，實際上硬體會將單次記憶體讀取操作廣播到半個執行緒束），廣播能夠在單個週期內發生，因此這個特性是非常有用的。雖然當所有16個執行緒都讀取相同地址時，這個功能可以極大提高效能，但當所有16個執行緒分別讀取不同的地址時，它實際上會降低效能。如果半個執行緒束中的所有16個執行緒需要訪問常量記憶體中的不同資料，那麼這個16次不同的讀取操作會被序列化，從而需要16倍的時間來發出請求。但如果從全域性記憶體中讀取，那麼這些請求就會同時發出。這種情況下，從常量記憶體讀取就會慢於從全域性記憶體中讀取。

需要注意的是，當我們宣告一個核心常量的時候，在編譯器將CUDA C程式碼轉換成PTX彙編程式碼時會用字面值（0x55555555）直接替換常量值(data)的地址。

const int data = 0x55555555;
int d = data;   //此時data會直接編譯為字面值0x55555555

但當我們宣告的是一個常量陣列時，編譯器在將C程式碼轉換成PTX彙編程式碼時將會使用陣列地址在彙編程式碼中。

const int data[3] = {0x11111111, 0x22222222, 0x33333333};
int d = data[1];    //此時data[1]會被編譯為data
 
[1]的地址

這時，在費米（計算能力為2.x的硬體）架構的裝置上，全域性記憶體藉助一級快取也能達到與常量記憶體相同的訪問速度。只有在計算能力為1.x的裝置上，由於全域性記憶體沒有用到快取技術，此時使用常量記憶體才會獲得明顯的效能提升。

2.常量記憶體的使用

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));

注意，在上例中，我們用cudaMemcpyToSymbol分別將資料從主機端複製到GPU的常量記憶體中和全域性記憶體中。此時，注意cudaMemcpyToSymbol函式的工作原理，該函式可以將資料複製到GPU上任何以全域性符號命名的記憶體區域，無論該符號是全域性記憶體還是常量記憶體。