在Compute Capability 2.0（Fermi）發佈之後，我想知道是否還有任何用於共享內存的用例。也就是說，什麼時候使用共享內存比讓L1在後臺執行它的魔法更好？CUDA：何時使用共享內存以及何時依賴L1緩存？
簡單的共享內存讓設計用於CC < 2.0的算法無需修改即可高效運行？
要通過共享內存進行協作，塊中的線程將寫入共享內存並與__syncthreads()同步。爲什麼不直接寫入全局內存（通過L1），並與__threadfence_block()同步？後一個選項應該更容易實現，因爲它不必涉及兩個不同的值的位置，並且它應該更快，因爲沒有從全局到共享內存的明確複製。由於數據被緩存在L1中，因此線程無需等待數據實際完成到全局內存。
使用共享內存時，可以保證在整個塊的持續時間內保留的值保留在那裏。這與L1中的價值觀是相反的，如果它們使用不夠頻繁，它們會被逐出。有沒有什麼情況下最好的辦法是將共享內存中很少使用的數據緩存起來，而不是讓L1根據算法實際使用的使用模式來管理它們？
就我所知，GPU中的L1緩存的行爲非常像CPU中的緩存。因此，你的評論「這與L1中的值相反，如果它們使用不夠頻繁就會被驅逐」對我來說沒有多大意義
當L1緩存中的數據不是被驅逐時足夠經常使用。通常，當請求一個以前不在緩存中的內存區域，並且其地址解析爲已經使用的內存區域時，它會被清除。我不知道NVidia採用的確切高速緩存算法，但假設常規的n路關聯，那麼每個內存條目只能緩存在整個高速緩存的一小部分中，基於它的地址
我想這個也可以回答你的問題。通過共享內存，您可以完全控制存儲在什麼地方的內容，而使用緩存時，所有內容都會自動完成。儘管編譯器和GPU在優化內存訪問方面仍然非常聰明，但有時您仍然可以找到更好的方法，因爲您知道輸入的內容以及哪些線程會做什麼程度當然）
即使數據在高速緩存中，但全局內存加載/存儲仍受制於合併規則，但共享內存在隨機訪問方面更加靈活。我嘗試使用L1緩存來存儲/計算直方圖，並且由於半隨機訪問模式，結果比使用共享內存慢得多。
此外，根據NVIDIA employee，當前L1高速緩存是直寫（立即寫入L2高速緩存），這會降低您的程序速度。
所以基本上，如果你真的想要性能的話，緩存會阻礙你。
計算能力2. *和3。*寫入時無效L1緩存線。計算能力3.0-3.5不會緩存L1中的全局讀取。在計算能力3. *設備上，每個存儲體8字節的共享存儲器帶寬實際上是256字節/時鐘，而L1被限制爲來自高速緩存行的128個字節。正如耶魯所說的共享內存有銀行衝突（所有的訪問必須是不同的銀行或同一個銀行的地址），而L1有地址分歧（所有地址必須在同一個128字節的高速緩存行），所以共享內存更有效率隨機訪問。 –
讓我對通用處理器（例如Intel AVX2有聚集，但它確實是串行）上實際上不存在SIMD內存訪問的原因提供一個猜想。我很確信這是因爲虛擬到物理地址轉換的成本很高，共享內存訪問不需要，因爲它是自己的地址空間。想象一下，必須並行執行32個TLB查找的代價！如果所有的32個地址都在同一個頁面中，可能會有一個優化？ –
謝謝，那確實回答了我的問題。我曾將緩存描繪爲能夠跟蹤哪些元素被最多使用，並且更願意緩存這些元素。我已經閱讀了關於n-way關聯緩存的內容，它看起來像主要的問題是它們可能會拋出一個經常使用的值，因爲另一個緩存線適合該插槽。 –
我認爲這意味着編寫CUDA程序的好策略通常可能是首先編寫算法以僅使用全局內存，然後查看L1是否工作良好，從而隱藏內存延遲。然後考慮手動優化與共享內存，如果算法結果是內存限制。 –