全局內存訪問優化（Coalesced Access）

什么是 Coalesced Access？

定義：一個 warp（32 個線程）在同一指令中訪問全局內存時，如果這些訪問請求可以合并成盡可能少的內存事務（通常是 32、64 或 128 字節對齊的塊），就叫 coalesced。

條件:一個 warp 的線程訪問 連續且對齊 的地址。

int tid = threadIdx.x + blockIdx.x * blockDim.x;
float val = d_array[tid];  // ? 連續訪問 → Coalesced

優化技巧

以結構體對齊內存：使用 __align__(16) 或 float4。

為什么需要對齊？

• GPU 內存總線要求訪問按 32/64/128 字節對齊，這樣才能合并成一次事務。

• 如果內存對齊不好，warp 訪問會拆分成多個事務，帶寬利用率降低。

技巧

• 使用 float4（4 個 float 一起）保證 16 字節對齊。

• 或者使用 CUDA 對齊修飾符：

float4 vs float

float

• 單個 4 字節（32 bit）的浮點數。

• 每個線程訪問 1 個 float 時，如果 warp 中 32 個線程訪問地址連續（0,1,2,3...），CUDA 會把它們合并成1~2 個內存事務，性能好。

float4

• CUDA 提供的 矢量類型，表示 4 個連續的 float（總共 16 字節）。

• 優點：

  • 天然 16 字節對齊（滿足 GPU 內存事務對齊要求）。

  • 每個線程一次加載 4 個浮點數，減少指令數，提高帶寬利用率。

//設置float4
float4 f;
f.x = 1.1, f.y = 2.2, f.z = 3.3, f.w = 4.4;
float4 v = data[idx];  // 讀取 4 個 float
result = v.x + v.y + v.z + v.w;

如果要讀取大數組，使用 float4 可以讓 每個線程批量讀取，提高 coalesced 訪問效率。

對比：

• 32 個線程一次訪問 float → 128 字節（32*4）

• 32 個線程一次訪問 float4 → 512 字節（32*16），如果對齊良好，GPU 可以用更少的事務完成。

__align__(n) 關鍵字

作用

• 強制結構體或變量的起始地址對齊到 n 字節邊界。

• 為什么？因為 GPU（和 CPU）要求數據按一定字節對齊訪問，否則：

  • 拆分訪問 → 多次內存事務 → 性能差

  • 未對齊訪問 → 有的設備直接報錯

struct __align__(16) MyStruct {float x, y, z, w;
}; // 占 16 字節，起始地址必須是 16 的倍數

• 如果不加 __align__(16)，可能被編譯器按 4 字節對齊排布，不符合 GPU 要求。

為什么 CUDA 推薦使用 float4 + 對齊？

• 全局內存的訪問規則：按 32/64/128 字節事務合并。

• 如果 warp 32 線程訪問 float（每個 4 字節），正好 128 字節，可以合并。

• 如果 warp 32 線程訪問 float4（每個 16 字節），正好 512 字節，GPU 需要 4 個事務，但每個事務更大，吞吐率更高。

• 重要：必須保證起始地址按 float4 對齊，否則性能下降。

行優先存儲，避免跨行訪問

• CUDA 全局內存是按一維線性存儲的，如果訪問跨行，會破壞 coalesced。

• 例如，二維矩陣 A[M][N]，默認按行優先（row-major）存儲：

內存布局: A[0][0], A[0][1], ..., A[0][N-1], A[1][0], ...

錯誤訪問模式（列遍歷）：

val = A[col][row]; // 每個線程跨 stride 訪問

每個線程 stride 大，warp 訪問不連續，性能差。

優化

• 保證 threadIdx.x 對應 最快變化維度（行訪問），這樣 warp 線程連續訪問。

調換索引順序，確保 threadIdx.x 是最快變化維度

• 原則：warp 線程訪問地址必須連續。

• 如果你的算法天然是列操作，可以調整線程分布：

如果是行優先，那么變化最快的其實是列下標，threadIdx.x對應的也應該是列。如果算法要求列優先，可以對row和col進行調換

以下是行優先情況下col和row的寫法

int col = blockIdx.x * blockDim.x + threadIdx.x;  // x 對應列
int row = blockIdx.y * blockDim.y + threadIdx.y;  // y 對應行

使用 Shared Memory 緩存 tile

為什么？

• 全局內存訪問延遲大（400~600 cycles），共享內存延遲低（≈100x 更快）。

• 如果每個線程直接從全局內存多次訪問，會拖慢性能。

• 解決：把要用的數據塊（tile）加載到共享內存，所有線程復用，減少全局訪問。

例子見上一篇文章：cuda編程筆記（9）--使用 Shared Memory 實現 tiled GEMM -CSDN博客

Bank Conflict

Bank Conflict（共享內存銀行沖突） 是 CUDA 編程中的一個性能問題，發生在多個線程同時訪問 共享內存（Shared Memory） 時。

共享內存的結構

• CUDA 的 共享內存被劃分成多個 Bank，類似一個并行訪問的“多路存儲器”。

• 每個 Bank 可以在一個時鐘周期內處理 1 個 32-bit 訪問請求。

• Warp（32 個線程）同時訪問共享內存時：

  • 如果 32 個線程訪問 32 個不同的 Bank → 無沖突（完美并行）。

  • 如果 多個線程訪問同一個 Bank 的不同地址 → 發生 Bank Conflict，訪問會被 串行化，性能大幅下降。

具體原理

假設：

• 共享內存被分為 32 個 Bank

• 每個 Bank 寬度 = 4 字節（一個 float）

• 地址映射公式：

bank_id = (address_in_bytes / 4) % 32

例子

__shared__ float s[32][32];
按行優先存儲（Row-major）：

• s[i][j] 的地址 = base + (i * 32 + j) * 4

情況 1：訪問同一列

如果每個線程訪問 s[threadIdx.x][k]（同一列 k），

• 地址 = base + (threadIdx.x * 32 + k) * 4

• bank_id = (threadIdx.x * 32 + k) % 32 = k（因為 threadIdx.x * 32 是 32 的倍數）

• 所有線程訪問同一 Bank（k） → 嚴重沖突

情況 2：訪問同一行

如果每個線程訪問 s[k][threadIdx.x]（同一行 k），

• 地址 = base + (k * 32 + threadIdx.x) * 4

• bank_id = (k * 32 + threadIdx.x) % 32 = threadIdx.x

• 每個線程訪問不同 Bank → 無沖突

避免 Bank Conflict 的方法

核心原則：讓 warp 內的 32 個線程訪問的地址盡量分布到不同的 bank。

• 按行訪問而非按列：

  • 推薦：s[threadIdx.y][threadIdx.x]（X 對應列，變化最快）

• 增加 padding（填充列）：

  • 如果二維數組導致 bank 沖突，可以在第二維加一個“dummy 列”，讓 stride ≠ 32：

__shared__ float s[TILE_SIZE][TILE_SIZE + 1];

使用結構化數據（float4）或 align：

• 一次加載多個元素，減少 warp 的 bank 競爭。

?