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，其主要作用有：

· An intra-block thread communication channel 線程間交流通道

· A program-managed cache for global memory data可編程cache

· Scratch pad memory for transforming data to improve global memory access patterns

本文主要涉及兩個例子作解釋：reduction kernel，matrix transpose kernel。

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數組才能這樣動態使用。

Memory Banks

為了獲得高帶寬，shared Memory被分成32（對應warp中的thread）個相等大小的內存塊，他們可以被同時訪問。不同的CC版本，shared memory以不同的模式映射到不同的塊（稍后詳解）。如果warp訪問shared Memory，對于每個bank只訪問不多于一個內存地址，那么只需要一次內存傳輸就可以了，否則需要多次傳輸，因此會降低內存帶寬的使用。

?

總結

以上是生活随笔為你收集整理的cuda的shared momery的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。