在cuda平行計算中，共享記憶體在gpu速度優化上扮演著重要作用，但是如果共享記憶體使用不當，也會導致速度不快反降或者提速效果不佳，如發生bank conflict;
bank的中文翻譯為儲存體，gpu 共享記憶體是基於儲存體切換的架構（bank-switched-architecture），一般現在的gpu都包含32個儲存體,即共享記憶體被分成了32個bank；根據gpu計算能力的不同（compute capability），每個共享記憶體儲存體的寬可以是32位（cc2.x)或64位(cc3.x以上),即連續的32-bits（或64-bits)字被分配到連續的32個bank中（計算能力不是描述gpu裝置計算能力強弱的絕對指標，他是相對的，準確的說他是乙個架構的版本號，他可以通過cudadevicesetsharedmemconfig() 配置成 cudasharedmembanksizefourbyte 四個位元組或者 cudasharedmembanksizeeightbyte（cc3.x以上) 。設定成8位元組可以有效避免雙精度資料的bank conflicts,預設是4位元組）, 但是這又遇到乙個問題，以telsa p100為例，我們切換bank的寬為32bit,即4個位元組，那麼32個bank僅僅為128b的記憶體，而telsa p100的共享記憶體為48kb,那麼多餘的記憶體呢？
我們看到這32bit我們定義為寬，那麼有寬就有高，在這個部落格中博主進行了這樣的比喻：
在共享記憶體中，連續的32-bits字被分配到連續的32個bank中，這就像電影院的座位一樣：一列的座位就相當於乙個bank，所以每行有32個座位，在每個座位上可以「坐」乙個32-bits的資料(或者多個小於32-bits的資料，如4個char型的資料，2個short型的資料)；而正常情況下，我們是按照先坐完一行再坐下一行的順序來坐座位的，在shared memory中位址對映的方式也是這樣的
其中0-31為bank編號,如果申請乙個共享記憶體陣列__shared__ int cache[64],int 恰好為4個位元組，那麼cache[0]訪問bank[0][0], cache[1]訪問bank[0][1],...,cache[31]訪問bank[0][31],cache超過32時，cache就會去訪問下一行的bank,即cache[32]就會訪問
bank[1][0],以此類推。
bank衝突就是在這樣的條件下產生，即如果乙個warp的多個執行緒訪問同乙個bank的不同字段時(注：不同欄位如bank[0][0],bank[1][0],...,bank[n][0])，那麼就發生了bank衝突，因為不同bank可以同時訪問，而當如果多個執行緒請求的記憶體位址被對映到了同乙個bank上，那麼這些請求就變成了序列的。
在bank conflicts中，我們一直在強調同一warp,這是因為warp是gpu執行時的排程單位，即對於gpu的每個sm執行的乙個block，事實上每一次僅有32個執行緒在同時執行，只是因為乙個gpu有多組sm，每個sm可以同時處理多個block，所以同時處理的執行緒數也就多了。因此不同warp訪問同一bank並不會造成衝突，因為事實上不同warp本來就不會同時訪問bank。
下面來看乙個bank conflict的例子：
__global__ void kernel1() //沒有bank conflict
__global__ void kernel2() //有bank conflict
int main()
可以看到，kernel1和kernel2的執行時間差別並不大，甚至kernel1還略大於kernel2，這是為什麼呢？
我查了很多資料，其中有一種說法，核函式啟動也是需要時間的，一般是us級別, 但是對於連續的核函式啟動，後面的kernel啟動延遲可以被隱藏掉（包括啟動隱藏和執行隱藏）
我們看到本例啟動的執行緒並不多，程式並不複雜，執行時間可能被隱藏時間抵消掉了，但是我們通過nvvp也可以看到效果（nvvp是nvprof的 圖形版）
通過nvvp gpu details分析，我們看到下圖：
(注：nvprof,nvvp和cuda提供的計時函式cudaeventrecord計算出來的時間都不一樣（有可能是硬體隨機化，但是差別有點大，不太像，此處還不清楚）
從shared memory efficiency處可以看到，kernel1的效率為100%，kernel2的效率為25%，剛好是我們前面分析的4倍。
前面我們定義bank conflict為乙個warp多個執行緒訪問同乙個bank的不同字段，那麼乙個warp多個執行緒訪問同乙個bank的相同字段
呢？如同時訪問bank[0][0]。結論是不會發生bank conflict,這就牽涉到gpu的廣播和多播機制，詳情可以檢視部落格
CUDA 共享記憶體
於gpu上啟動的每個執行緒塊上的共享記憶體，cuda c編譯器都會建立該變數的乙個副本。同一執行緒塊的每個執行緒都共享這塊記憶體，但是執行緒無法看到也不能修改其他執行緒塊中的共享記憶體。這樣做可以使得乙個執行緒塊中的多個執行緒能夠在計算上通訊和協作。共享記憶體緩衝區駐留在物理gup上，因此訪問共享記...
共享CUDA記憶體
共享cuda記憶體 程序間共享 此功能僅限於linux。將裝置陣列匯出到另乙個程序 使用cuda ipc api，可以與同一臺計算機上的另乙個程序共享裝置陣列。為此，請使用.get ipc handle 裝置陣列上的方法獲取乙個ipcarrayhandle物件，該物件可以轉移到另乙個程序。devic...
CUDA申請動態共享記憶體
直接上 include include global void kernel int main return 0 程式輸出 1021.000000 1022.000000 1023.000000 index 1021.000000 index 1022.000000 index 1023.00000...