SIMT 記憶體效能 (SIMT Memory Performance)

#cuda #memory #memory-coalescing #bank-conflict #shared-memory #concept

正確使用記憶體是寫出高效能 CUDA kernel 的關鍵。本筆記用一個「矩陣轉置 (matrix transpose)」範例，循序漸進地改善 global memory 與 shared memory 的存取，帶出兩個核心效能主題：coalescing（合併）與 bank conflict（bank 衝突）。

重點總覽

項目	重點
Global memory transaction	global memory 以 32-byte transaction 為單位存取；一個 warp 把 32 個 thread 的請求合併成必要數量的 transaction
Coalesced access	warp 內 threads 存取同一批 32-byte segment（連續或排列皆可）；最佳化目標 = 最大化「used bytes / transferred bytes」
完美合併	連續 thread 取連續 4-byte word → 128 bytes → 4 個 32-byte transaction → 100% 利用率
病態未合併	連續 thread 取相距 ≥32 bytes 的元素 → 每 thread 1 個 transaction → 32×32=1024 bytes 流量、僅用 128 bytes → 12.5% 利用率
Transpose (global)	讀 `a` 合併、寫 `c` 不合併（stride = ld）；ld>32 時即病態案例
Shared memory banks	shared memory 分 32 個 bank，連續 32-bit word 對映到連續 bank，每 bank 頻寬 32 bits/clock
Bank conflict	同 warp 多 thread 存取同一 bank 的不同位置 → 序列化 → 效能損失
例外（無衝突）	同 warp 多 thread 存取同一位置：讀 → broadcast；寫 → 僅一 thread 寫入（哪個未定義）
Transpose (shared)	用 shared memory 當 user-managed cache 暫存 → 讀寫 global memory 皆合併；需 `__syncthreads()`
Padding 修正	把 column 維度 +1（`[32][33]`）即可消除轉置時的 32-way bank conflict

2.3.4.1 Coalesced Global Memory Access（global memory 合併存取）

Global memory 以 32-byte memory transaction 為單位存取。 當一個 CUDA thread 向 global memory 請求一個 word，所屬 warp 會把該 warp 內所有 threads 的請求 coalesce（合併） 成「滿足這些請求所需的最少 transaction 數」，數量取決於每個 thread 存取的 word 大小以及這些位址在 threads 間的分佈。

一個 thread 請求 4-byte word，warp 實際送出的 transaction 仍是 32 bytes。
要最有效率使用記憶體系統，warp 應用滿一次 transaction 取回的所有資料；理想是讓同 warp 其他 threads 也用到那 32 bytes 裡的其他 4-byte word。

完美合併（best case）：warp 內連續 threads 請求記憶體中連續的 4-byte word：

連續 thread 取連續 4-byte word：32 threads × 4 bytes = 128 bytes
被切成 4 個 32-byte transaction，每個都被 8 個 thread 完整用滿

addr →  [ 0..31 ][ 32..63 ][ 64..95 ][ 96..127 ]   (bytes)
thread   T0..T7   T8..T15   T16..T23  T24..T31
txn      ── 1 ──  ── 2 ──   ── 3 ──   ── 4 ──
利用率 = 128 used / 128 transferred = 100%

病態未合併（worst case）：連續 threads 存取彼此相距 ≥32 bytes 的元素，warp 被迫為每個 thread 各發一個 32-byte transaction：

每個 thread 落在不同 32-byte segment（stride ≥ 32 bytes）
T0 → [seg A 32B] 只用 4B
T1 → [seg B 32B] 只用 4B
...
T31 → [seg ... 32B] 只用 4B
流量 = 32 bytes × 32 threads = 1024 bytes
實際使用 = 4 bytes × 32 = 128 bytes
利用率 = 128 / 1024 = 12.5%   ← 非常沒效率

等價觀點（用 transaction 數思考）：一個 warp 的單一 load 指令會請求 32 個位址，需要幾個 global memory transaction 來滿足？

情況	所需 transaction 數
最佳（理論極限）	1 個 transaction 滿足全部
4-byte 資料完美合併	4 個 transaction
最差	最多 32 個 transaction

一般而言，滿足一個 load 所需的 transaction 數越少，效能越好。
不要求連續 threads 必須存取連續記憶體：只要 warp 內不同 threads 以線性或排列 (permuted) 方式存取同一批 32-byte segment，就算合併。最佳化原則就是 最大化 used bytes / transferred bytes 的比值。

vecAdd 是天然合併的例子：連續 workIndex（= 連續 threads）存取陣列連續元素。

__global__ void vecAdd(float* A, float* B, float* C, int vectorLength)
{
    int workIndex = threadIdx.x + blockIdx.x * blockDim.x;
    if (workIndex < vectorLength)
        C[workIndex] = A[workIndex] + B[workIndex];
}

每個 thread 對三個陣列的存取都隨 threadIdx.x 連續推進 → 完美合併。

Important

確保 global memory 存取正確合併，是寫出高效能 CUDA kernel 最重要的效能考量之一。應用程式必須盡可能有效率地使用記憶體系統。

2.3.4.1.1 Matrix Transpose Using Global Memory（global memory 轉置的非合併問題）

一個 out-of-place 矩陣轉置：把 N×N 的 32-bit float 方陣由 a 轉到 c。用 2d grid、32×32 的 2d thread block（blockDim.x = blockDim.y = 32），每個 block 處理一個 32×32 的 tile，每個 thread 處理一個唯一元素，故不需 thread 同步。

/* 以 2D 索引存取 row-major 1D 陣列；ld = leading dimension = 矩陣的欄數 */
#define INDX(row, col, ld) (((row) * (ld)) + (col))

__global__ void cuda_transpose(int m, float *a, float *c)
{
    int myCol = blockDim.x * blockIdx.x + threadIdx.x;
    int myRow = blockDim.y * blockIdx.y + threadIdx.y;
    if (myRow < m && myCol < m)
        c[INDX(myCol, myRow, m)] = a[INDX(myRow, myCol, m)];  // 寫 c 不合併
}

判斷是否合併，要看「連續 threads 是否存取連續記憶體」。在 2d block 中 threadIdx.x 變化最快，故連續 threadIdx.x（出現在 myCol）應對映連續記憶體：

存取	`myCol`(=threadIdx.x) 在 INDX 的位置	隨 threadIdx.x++ 位址變化	結果
讀 `a[INDX(myRow, myCol, m)]`	第二個參數 (col)	+1 元素（連續）	完美合併
寫 `c[INDX(myCol, myRow, m)]`	第一個參數 (row)	+ld 元素（stride=ld）	不合併

Warning

寫 c 時，連續 threadIdx.x 寫入相距 ld 個元素的位置。只要矩陣大於 32（即 ld > 32），stride 就 ≥32 個 float = 128 bytes ≥ 32 bytes，正是 Figure 13 的病態未合併案例。解法：引入 shared memory 暫存。

2.3.4.2 Shared Memory Access Patterns（shared memory 存取樣式與 bank）

Shared memory 分成 32 個 bank，組織方式為「連續的 32-bit word 對映到連續的 bank」，每個 bank 頻寬為 32 bits / clock cycle。

32-bit word 索引:  0   1   2   3  ... 31  32  33 ...
bank 編號:         0   1   2   3  ... 31   0   1 ...   (word % 32)

Bank conflict：同一 warp 內多個 threads 嘗試存取同一個 bank 中的不同元素時發生。該 bank 的存取會被序列化 (serialized)，直到所有請求該 bank 的 threads 都拿到資料 → 造成效能損失。
理想情境：warp 內每個 thread 存取的 shared memory 位置都落在不同 bank（不同顏色）。

Tip

兩個不會造成衝突的例外 —— 同一 warp 多 threads 存取同一個 shared memory 位置時：

讀取：該 word 被 broadcast 給所有請求的 threads。
寫入：每個 shared memory 位址只由其中一個 thread 寫入（哪個 thread 寫入是 undefined）。

strided 存取（以 32-bit bank size 為例）：

stride	結果
1 個 32-bit word	無 bank conflict（每 thread 落在不同 bank）
2 個 32-bit word	two-way bank conflict
3 個 32-bit word	無 bank conflict

Important

避免 bank conflict 是寫出使用 shared memory 的高效能 kernel 的重要效能考量。

2.3.4.2.1 Matrix Transpose Using Shared Memory（用 shared memory 改善合併）

把 shared memory 當成 user-managed cache，暫存 (stage) 來自 global memory 的 load/store，使讀與寫 global memory 都合併。注意此處 INDX 改為 column-major（(col)*(ld)+(row)，ld = 列數）。

#define THREADS_PER_BLOCK_X 32
#define THREADS_PER_BLOCK_Y 32
/* column-major 索引；ld = leading dimension = 矩陣的列數 */
#define INDX(row, col, ld) (((col) * (ld)) + (row))

__global__ void smem_transpose(int m, float *a, float *c)
{
    __shared__ float smemArray[THREADS_PER_BLOCK_X][THREADS_PER_BLOCK_Y];
    const int tileRow = blockDim.x * blockIdx.x;   // 此 block 的 tile 起點
    const int tileCol = blockDim.y * blockIdx.y;

    /* 1) 從 global memory 讀進 shared memory（讀 a 合併） */
    smemArray[threadIdx.x][threadIdx.y] =
        a[INDX(tileRow + threadIdx.y, tileCol + threadIdx.x, m)];

    __syncthreads();   // 確保整個 block 都載入完成才能讀 smem

    /* 2) 由 shared memory 寫回 global memory（寫 c 也合併） */
    c[INDX(tileCol + threadIdx.y, tileRow + threadIdx.x, m)] =
        smemArray[threadIdx.y][threadIdx.x];   // 讀 smem 時交換 x/y → 完成轉置
}

關鍵步驟與原理：

計算 tile 起點：tileRow/tileCol 由 block 索引決定；同一 block 內所有 threads 共用同一組 tile 索引（tile 起始位置）。
載入到 shared memory：每個 block 複製 32×32 tile 到 smem，由 32 個 warp 執行、warp 間無保證順序。threadIdx.x 位於 INDX 的第二個參數 → 連續 threads 取連續 a → 讀 a 完美合併。
__syncthreads()：保證 block 內所有 threads 都完成寫入 smem，下一步才能安全地從 smem 讀。

Warning

缺少 __syncthreads() 時，無法保證所有 warp 都已把 a 寫進 smem，部分 warp 可能提前去讀 smem → 讀到未載入的資料。只要 thread 處理／儲存它自己沒有載入的資料，就必須同步以確保該元素的 load 已完成。

轉置與寫回：讀 smemArray 時交換 threadIdx.x／threadIdx.y（轉置 tile 內元素）；寫 c 時交換 tileRow／tileCol（把 tile 放到正確位置）；並讓 threadIdx.x 留在 INDX 第二參數 → 寫 c 也合併。

@cuda.jit
def smem_transpose(a, c):
    smemArray = cuda.shared.array(shape=(32, 32), dtype=np.float32)
    tile_col = cuda.blockDim.x * cuda.blockIdx.x
    tile_row = cuda.blockDim.y * cuda.blockIdx.y
    smemArray[cuda.threadIdx.x, cuda.threadIdx.y] = \
        a[tile_row + cuda.threadIdx.y, tile_col + cuda.threadIdx.x]
    cuda.syncthreads()
    c[tile_col + cuda.threadIdx.y, tile_row + cuda.threadIdx.x] = \
        smemArray[cuda.threadIdx.y, cuda.threadIdx.x]

Tip

此 kernel 展示 shared memory 的兩個常見用途：

Stage（暫存） global memory 資料，讓讀與寫 global memory 都合併。
共享資料：讓同一 block 內的 threads 彼此分享資料。

2.3.4.2.2 Shared Memory Bank Conflicts（bank 衝突成因與 padding 解法）

上一個 shared memory 轉置 kernel 已做到 global memory 合併，但沒有檢查 shared memory bank conflict。考慮宣告：

__shared__ float smemArray[32][32];

以 32×32 的 2d block 啟動，warp 內每個 thread 的 threadIdx.y 固定、0 <= threadIdx.x < 32。由於 C++ 與 Python 多維陣列都是「最後一個索引變化最快」：

smemArray[row][col]，col 變化最快 → 線性位址 = row*32 + col
bank = (row*32 + col) % 32 = col % 32   （與 row 無關）

按「整欄」存取 (col 固定, row=0..31)：       按「整列」存取 (row 固定, col=0..31)：
 smemArray[0][0], [1][0] ... [31][0]         smemArray[0][0], [0][1] ... [0][31]
 stride = 32 元素 → 全部落在 bank 0          stride = 1 元素 → 落在 bank 0..31
 → 32-way bank conflict (序列化)             → 無 bank conflict
   bank: 0 0 0 0 ... 0                          bank: 0 1 2 3 ... 31

回頭檢查轉置 kernel 的兩處 shared memory 存取：

存取	程式碼	`threadIdx.x` 是第幾索引	smem stride	結果
寫入 smem	`smemArray[threadIdx.x][threadIdx.y]`	第一索引	32 元素	32-way bank conflict
讀取 smem	`smemArray[threadIdx.y][threadIdx.x]`	第二索引	1 元素	無 bank conflict

Warning

因為陣列是 row-major，threadIdx.x 當第一個索引時，連續 threads 以 stride 32 存取 smem → 全落同 bank → 32-way conflict。此 kernel 因此有「一處無衝突、一處 32-way 衝突」。

Padding 修正：把 column 維度 +1，使列長變 33（不是 2 的次方），就能讓「整欄」存取的位址錯開到不同 bank：

__shared__ float smemArray[THREADS_PER_BLOCK_X][THREADS_PER_BLOCK_Y + 1];  // [32][33]

smemArray = cuda.shared.array(shape=(32, 32 + 1), dtype=np.float32)  # [32][33]

padding 後 smemArray[32][33]，整欄存取 (col 固定, row=0..31)：
線性位址 = row*33 + col → bank = (row*33 + col) % 32 = (row + col) % 32
row=0..31 → bank 依序 +1 遞增，落在 32 個不同 bank → 無 conflict

無論「整欄」或「整列」存取，warp 內 threads 都落在不同 bank（不同顏色）→ 衝突消除

Tip

Padding 之所以有效：列長由 32（= bank 數）改成 33，使「同一欄、不同列」的元素之間相差 33（與 32 互質），對 32 取餘數後恰好落在全部不同的 bank。代價是每列多一個未用的 float 空間。

考試/測驗重點

情境/關鍵字	答案
global memory transaction 大小	32 bytes
一個 thread 請求 4-byte word，實際 transaction 多大	32 bytes
warp 完美合併存取 4-byte 資料需幾個 transaction	4 個（128 bytes / 32 bytes）
完美合併利用率	100%
病態最差利用率	12.5%（128 used / 1024 transferred）
最差情況流量	32 bytes × 32 threads = 1024 bytes
合併的真正條件	warp 內 threads 存取同一批 32-byte segment（連續或 permuted 皆可），最大化 used/transferred
連續 threads 必須存取連續記憶體才能合併？	否（只是最常見的達成方式）
2d block 哪個維度變化最快	`threadIdx.x`
global 轉置：讀 a / 寫 c 是否合併	讀 a 合併、寫 c 不合併（stride = ld）
寫 c 何時變成病態案例	ld > 32（矩陣大於 32）時，stride ≥ 32 bytes
shared memory 有幾個 bank	32 個，每 bank 32 bits/clock
bank 對映規則	連續 32-bit word → 連續 bank（`word % 32`）
bank conflict 定義	同 warp 多 threads 存取同一 bank 的不同位置 → 序列化
同 warp 存取同一位置（讀）	broadcast，無衝突
同 warp 存取同一位置（寫）	僅一 thread 寫入，哪個 undefined，無衝突
stride 2 word vs stride 3 word	stride 2 → two-way conflict；stride 3 → 無衝突
`smemArray[32][32]` 按整欄存取	32-way bank conflict（stride 32）
`smemArray[32][32]` 按整列存取	無衝突（stride 1）
轉置 kernel 哪一處有衝突	寫入 smem `smemArray[threadIdx.x][threadIdx.y]`（threadIdx.x 是第一索引）→ 32-way
消除 bank conflict 的常見手法	padding：column 維度 +1，宣告成 `[32][33]`
為何寫不載入的資料需要同步	`__syncthreads()` 保證 load 完成，避免讀到未寫入的 smem
shared memory 兩大用途	(1) stage 資料使 global 存取合併 (2) block 內 threads 共享資料

重點總覽

2.3.4.1 Coalesced Global Memory Access（global memory 合併存取）

2.3.4.1.1 Matrix Transpose Using Global Memory（global memory 轉置的非合併問題）

2.3.4.2 Shared Memory Access Patterns（shared memory 存取樣式與 bank）

2.3.4.2.1 Matrix Transpose Using Shared Memory（用 shared memory 改善合併）

2.3.4.2.2 Shared Memory Bank Conflicts（bank 衝突成因與 padding 解法）

考試/測驗重點

Related Notes