CUDA SHARED MEMORY
在global Memory部分，資料對齊和連續是很重要的話題，當使用L1的時候，對齊問題可以忽略，但是非連續的獲取記憶體依然會降低效能。依賴於演算法本質，某些情況下，非連續訪問是不可避免的。使用shared memory是另一種提高效能的方式。
GPU上的memory有兩種：
-
On-board memory
-
On-chip memory
global memory就是一塊很大的on-board memory，並且有很高的latency。而shared memory正好相反，是一塊很小，低延遲的on-chip memory，比global memory擁有高得多的頻寬。我們可以把他當做可程式設計的cache，其主要作用有：
shared memory（SMEM）是GPU的重要組成之一。物理上，每個SM包含一個當前正在執行的block中所有thread共享的低延遲的記憶體池。SMEM使得同一個block中的thread能夠相互合作，重用on-chip資料，並且能夠顯著減少kernel需要的global memory頻寬。由於APP可以直接顯式的操作SMEM的內容，所以又被稱為可程式設計快取。
由於shared memory和L1要比L2和global memory更接近SM，shared memory的延遲比global memory低20到30倍，頻寬大約高10倍。
當一個block開始執行時，GPU會分配其一定數量的shared memory，這個shared memory的地址空間會由block中的所有thread 共享。shared memory是劃分給SM中駐留的所有block的，也是GPU的稀缺資源。所以，使用越多的shared memory，能夠並行的active就越少。
關於Program-Managed Cache：在C語言程式設計裡，迴圈（loop transformation）一般都使用cache來優化。在迴圈遍歷的時候使用重新排列的迭代順序可以很好利用cache區域性性。在演算法層面上，我們需要手動調節迴圈來達到令人滿意的空間區域性性，同時還要考慮cache size。cache對於程式設計師來說是透明的，編譯器會處理所有的資料移動，我們沒有能力控制cache的行為。shared memory則是一個可程式設計可操作的cache，程式設計師可以完全控制其行為。
Shared Memory Allocation
我們可以動態或者靜態的分配shared Memory，其宣告即可以在kernel內部也可以作為全域性變數。
其識別符號為：__ shared__。
下面這句話靜態的聲明瞭一個2D的浮點型陣列：
__ shared __ float tile[size_y][size_x];
如果在kernel中宣告的話，其作用域就是kernel內，否則是對所有kernel有效。如果shared Memory的大小在編譯器未知的話，可以使用extern關鍵字修飾，例如下面宣告一個未知大小的1D陣列：
extern shared int tile[];
由於其大小在編譯器未知，我們需要在每個kernel呼叫時，動態的分配其shared memory，也就是最開始提及的第三個引數：
kernel<<<grid, block, isize * sizeof(int)>>>(…)
應該注意到，只有1D陣列才能這樣動態使用。
Shared Memory Banks and Access Mode
之前博文對latency和bandwidth有了充足的研究，而shared memory能夠用來隱藏由於latency和bandwidth對效能的影響。下面將解釋shared memory的組織方式，以便研究其對效能的影響。
Memory Banks
為了獲得高頻寬，shared Memory被分成32（對應warp中的thread）個相等大小的記憶體塊，他們可以被同時訪問。不同的CC版本，shared memory以不同的模式對映到不同的塊（稍後詳解）。如果warp訪問shared Memory，對於每個bank只訪問不多於一個記憶體地址，那麼只需要一次記憶體傳輸就可以了，否則需要多次傳輸，因此會降低記憶體頻寬的使用。
Bank Conflict
當多個地址請求落在同一個bank中就會發生bank conflict，從而導致請求多次執行。硬體會把這類請求分散到儘可能多的沒有conflict的那些傳輸操作 裡面，降低有效頻寬的因素是被分散到的傳輸操作個數。
warp有三種典型的獲取shared memory的模式：
-
Parallel access：多個地址分散在多個bank。
-
Serial access：多個地址落在同一個bank。
-
Broadcast access：一個地址讀操作落在一個bank。
Parallel access是最通常的模式，這個模式一般暗示，一些（也可能是全部）地址請求能夠被一次傳輸解決。理想情況是，獲取無conflict的shared memory的時，每個地址都在落在不同的bank中。
Serial access是最壞的模式，如果warp中的32個thread都訪問了同一個bank中的不同位置，那就是32次單獨的請求，而不是同時訪問了。
Broadcast access也是隻執行一次傳輸，然後傳輸結果會廣播給所有發出請求的thread。這樣的話就會導致頻寬利用率低。
下圖是最優情況的訪問圖示：
下圖一種隨機訪問，同樣沒有conflict：
下圖則是某些thread訪問到同一個bank的情況，這種情況有兩種行為：
-
Conflict-free broadcast access if threads access the same address within a bank
-
Bank conflict access if threads access different addresses within a bank
Synchronization
因為shared Memory可以被同一個block中的不同的thread同時訪問，當同一個地址的值被多個thread修改就導致了inter-thread conflict，所以我們需要同步操作。CUDA提供了兩類block內部的同步操作，即：
· Barriers
· Memory fences
對於barrier，所有thread會等待其他thread到達barrier point；對於Memory fence，所有thread會阻塞到所有修改Memory的操作對其他thread可見，下面解釋下CUDA需要同步的主要原因：weakly-ordered。
Weakly-Ordered Memory Model
現代記憶體架構有非常寬鬆的記憶體模式，也就是意味著，Memory的獲取不必按照程式中的順序來執行。CUDA採用了一種叫做weakly-ordered Memory model來獲取更激進的編譯器優化。
GPU thread寫資料到不同的Memory的順序（比如shared Memory，global Memory，page-locked host memory或者另一個device上的Memory）同樣沒必要跟程式裡面順序呢相同。一個thread的讀操作的順序對其他thread可見時也可能與實際上執行寫操作的thread順序不一致。
為了顯式的強制程式以一個確切的順序執行，就需要用到fence和barrier。他們也是唯一能保證kernel對Memory有正確的行為的操作。
Explicit Barrier
同步操作在我們之前的文章中也提到過不少，比如下面這個：
void __syncthreads();
__syncthreads就是作為一個barrier point起作用，block中的thread必須等待所有thread都到達這個point後才能繼續下一步。這也保證了所有在這個point之前獲取global Memory和shared Memory的操作對同一個block中所有thread可見。__syncthreads被用來協作同一個block中的thread。當一些thread獲取Memory相同的地址時，就會導致潛在的問題（讀後寫，寫後讀，寫後寫）從而引起未定義行為狀態，此時就可以使用__syncthreads來避免這種情況。
使用__syncthreads要相當小心，只有在所有thread都會到達這個point時才可以呼叫這個同步，顯而易見，如果同一個block中的某些thread永遠都到達該點，那麼程式將一直等下去，下面程式碼就是一種錯誤的使用方式：
if (threadID % 2 == 0) {
__syncthreads();
} else {
__syncthreads();
}
Memory Fence
這種方式保證了任何在fence之前的Memory寫操作對fence之後thread都可見，也就是，fence之前寫完了，fence之後其它thread就都知道這塊Memory寫後的值了。fence的設定範圍比較廣，分為：block，grid和system。
可以通過下面的API來設定fence：
void __threadfence_block();
看名字就知道，這個函式是對應的block範圍，也就是保證同一個block中thread在fence之前寫完的值對block中其它的thread可見，不同於barrier，該function不需要所有的thread都執行。
下面是grid範圍的API，作用同理block範圍，把上面的block換成grid就是了：
void __threadfence();
下面是system的，其範圍針對整個系統，包括device和host：
void __threadfence_system();