CUDA 共享記憶體深度解析：Bank Conflict、同步與效能優化實戰

CUDA 共享記憶體深度解析：Bank Conflict、同步與效能優化實戰

共享記憶體是 GPU 片上可編程快取，延遲比全域記憶體低 20-30 倍、頻寬高 10 倍。善用它需掌握 Bank 組織、存取模式、同步原語、資料佈局。

硬體架構

• 每 SM 擁有 48KB/96KB/164KB 共享記憶體 (視架構)。
• 劃分為 32 個 Bank (對應 Warp 32 執行緒)，可並行存取。
• Fermi (CC 2.x)：4 Byte/Bank，bank = (addr/4)%32。
• Kepler+ (CC 3.x+)：8 Byte/Bank，bank = (addr/8)%32，相鄰 64-bit 同 Bank 無衝突。
• 可動態切換：cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeFourByte/EightByte)。

三種存取模式與衝突

關鍵優化技巧

1. 二維陣列佈局：Row-Major 存取

__shared__ float tile[32][32];
// 好：連續 threadIdx.x 存取連續地址 → 無衝突
val = tile[threadIdx.y][threadIdx.x];
// 差：threadIdx.x 變化導致同 Bank → 32-way 衝突 (Fermi)
val = tile[threadIdx.x][threadIdx.y];

2. Padding 消除 Bank Conflict

// Fermi: +1 column 即可
__shared__ float tile[32][33]; // 32*33*4=4224B
// Kepler (8B/Bank): 需根據陣列大小試誤，可能需 +2/+4

3. 動態共享記憶體 + 手動索引

extern __shared__ int tile[]; // 1D
int row_idx = threadIdx.y * blockDim.x + threadIdx.x;
int col_idx = threadIdx.x * blockDim.y + threadIdx.y;
tile[row_idx] = ...; // Row-major 寫
__syncthreads();
val = tile[col_idx]; // Column-major 讀 → 有衝突

4. 同步原語

• __syncthreads()：Block 級 Barrier，所有執行緒到達才繼續。必須所有執行緒均遇到，否則死鎖。
• __threadfence_block()：Block 級 Memory Fence，確保寫入對 Block 內其他執行緒可見。
• __threadfence() / __threadfence_system()：Grid/系統級 Fence。
• volatile：強制每次從共享/全域記憶體讀取，防編譯器快取優化。

實戰案例：歸約與矩陣轉置

• Reduction：樹狀歸約在共享記憶體中，每步 __syncthreads()，避免 Warp Divergence。
• Matrix Transpose：以 32x32 Tile 為單位，Row-major 讀全域 → 共享記憶體 → Column-major 寫全域，配合 Padding 消除共享記憶體衝突。

效能分析工具

nvprof --metrics shared_load_transactions_per_request,shared_store_transactions_per_request ./kernel
// 理想值 = 1.0，>1 表示有 Bank Conflict

1. 上一篇文章： C 語言入門筆記（一）：計算機基礎、語言發展史與 C 語言特性
2. 下一篇文章： 有限狀態機編程模式：原理、實現與應用場景

關於C語言

C語言動態內存分配：malloc、calloc 與 realloc 的區別、聯系與使用陷阱

函數原型 void* malloc(size_t size)

經典算法研究：手把手教你用C語言從零實現SIFT算法（上）

作者：July | 日期：2011年3月12日 | 參考：R

C# 三種複雜變數型別詳解

前面介紹變數都有基本取值範圍，除了 String 型別外，每

2048 遊戲 C 語言終端機版完整原始碼與邏輯解析

經典 2048 遊戲：4×4 終端機純 C 實作 2048