一、概述

       常量記憶體其實只是全域性記憶體的一種虛擬地址形式，並沒有特殊保留的常量記憶體塊。常量記憶體有兩個特徵，一個是快取記憶體，另一個是它支援將單個值廣播到執行緒束中的每個執行緒。

       常量記憶體，通過名字我們就能猜到它是隻讀記憶體。這種型別的記憶體要麼是在編譯時宣告為只讀記憶體，要麼是在執行時通過主機端定義為只讀記憶體。常量只是從GPU記憶體的角度而言。常量記憶體大小被限制為64K。

      在編譯時宣告一塊常量記憶體，需要用到__constant__ 關鍵字。

      如果要在執行時改變常量記憶體區中的內容，只需在呼叫GPU核心之前簡單地呼叫cudaCopyToSymbol 函式。如果在編譯階段或主機端執行階段沒有定義常量記憶體，那麼常量記憶體區將未定義。

二、常量記憶體快取記憶體

1、在計算能力為1.x 的裝置

     在計算能力為1.x 的裝置上，常量記憶體有一個特性，就是將資料快取到一塊8KB 的一級快取上，使隨後的訪問變得更加快速。但前提是記憶體中的資料可能在程式中重複利用。此外，廣播訪問也被高度優化，使得訪問相同記憶體地址的執行緒能夠在單個期間內完成。

2、在計算能力為2.x 的裝置

     在費米架構的裝置以及後續釋出的裝置上出現了二級快取。費米架構裝置上的二級快取可以在SM之間進行共享。所有的記憶體訪問將自動地快取到二級快取中。對任何宣告為常量的記憶體區進行基於非執行緒的訪問都將通過常量快取。所謂的基於非執行緒的訪問，即對所訪問的陣列的索引進行計算時不需要用到 threadIdx.x。

三、常量記憶體廣播機制

      常量記憶體有一個非常有用的特性，該特性主要將資料分配或廣播到執行緒束的每個執行緒中。廣播能在單個週期內發生，因此，該特性非常有用。

      在基於二級快取訪問機制的費米裝置上，同樣可以使用廣播機制，通過廣播機制，我們可以快速的將資料分配到執行緒束的多個執行緒中。每個執行緒從常量記憶體讀取元素N，這樣會觸發廣播機制，在讀到資料之後廣播到執行緒束中的每個執行緒。由於常量記憶體區幾乎可以到達與一級快取相同的速度，因此這類演算法使用常量記憶體效能非常好。

      然而，需要特別注意的是，如果一個常量只是字面值，那麼最好用 #define 對字面值進行定義，因為這樣可以減少常量記憶體的使用。

四、執行時進行常量記憶體更新

       GPU 上的常量記憶體並不是真正意義上的常量記憶體，因為GPU上並沒有專門常量記憶體預留特殊記憶體區。由於常量記憶體是通過16位地址進行訪問的，而16位地址能夠快速進行訪問，因此常量記憶體最大限制為64KB。這樣做會帶來一定好處但也會帶來一些問題。首先，通過調研cudaMemcpyToSymbol 函式，常量記憶體可以按塊或片的形式進行更新，一次最多能更新64K。

      注意cudaMemcpyToSymbol 函式的工作原理。該函式可以將資料複製到GPU上任何以全域性符合命名的記憶體區域，無論該符號是全域性記憶體還是常量記憶體。因此，我們可以將一塊64K大小的資料複製到一個64K 大小記憶體區上，從而通過常量記憶體快取進行訪問。當所有執行緒訪問同一資料元素時，這種訪問方式非常有用，因為我們可以藉助快取技術從常量記憶體區獲取然後在廣播到每個執行緒中。