CUDA全局內存訪問模式優化完全指南

CUDA全局內存訪問模式優化完全指南

GPU應用性能往往受限於內存帶寬，優化全局內存訪問是獲取高性能的首要任務。本文系統梳理對齊訪問、合併訪問、緩存層級、讀寫模式、數據佈局（AoS vs SoA）、循環展開等核心優化技術，並配合實測數據驗證。

一、內存層級與訪問單元

• 全局內存訪問經過L2緩存，部分架構經L1緩存
• L1緩存行：128 bytes，映射設備內存128字節對齊塊
• L2緩存段：32 bytes
• Warp（32線程）每線程請求4 bytes → 總計128 bytes，完美契合L1緩存行

二、核心優化目標：對齊與合併

理想模式：連續且對齊 → 單次128B傳輸，總線利用率100%

最差模式：離散非對齊 → 多次傳輸，大量無效數據，利用率極低

三、緩存模式下的讀操作分析

L1啟用（默認）

• 事務單元：128 bytes
• 連續對齊：1事務，100%效率
• 對齊但線程ID非連續：仍1事務（數據在同一緩存行）
• 非連續非對齊：2事務，50%效率
• 完全亂序：最多32事務，效率極低

注意：GPU L1僅針對空間局部性優化，無時間局部性，頻繁訪問不會駐留數據。

L1禁用（-Xptxas -dlcm=cg），僅走L2

• 事務單元：32 bytes
• 非對齊情況下無效傳輸大幅減少，效率從50%提升至80%
• 高佔用率場景下未對齊反而比L1啟用時表現更好

四、只讀緩存

• CC 3.5+可替代L1，單元32 bytes，通常表現優於L1
• 使用方式：__ldg(&ptr)內建函數 或 const __restrict__修飾指針

五、寫操作特性

• 不經L1，僅緩存於L2，單元32 bytes
• 事務可為1/2/4段（32/64/128 bytes）
• 理想：連續對齊 → 單次4段事務
• 離散：多次1段事務
• 對齊且64B範圍內 → 單次2段事務

六、實測驗證：偏移訪問影響

__global__ void readOffset(float *A, float *B, float *C, int n, int offset) {
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int k = i + offset;
    if (k < n) C[i] = A[k] + B[k];
}

未對齊導致L1事務數翻倍，禁用L1後事務數恢復正常但延遲略增。

七、數據佈局：AoS vs SoA

// AoS: 結構體數組
struct InnerStruct { float x, y; };
InnerStruct myAoS[N]; // x,y交替存儲

// SoA: 數組結構體
struct InnerArray { float x[N], y[N]; }; // 所有x連續，所有y連續

實測對比（Fermi M2070，1M元素）：

SoA預先連續佈局，完美滿足合併訪問，性能提升43%。

八、性能調優技巧

1. 循環展開

每線程處理4個元素，增加獨立內存操作並行度：

__global__ void readOffsetUnroll4(...) {
    unsigned int i = blockIdx.x * blockDim.x * 4 + threadIdx.x;
    // 處理 i, i+blockDim.x, i+2*blockDim.x, i+3*blockDim.x
}

實測：展開後耗時從1.8ms降至0.6ms（3倍提升），對齊/未對齊差異縮小但仍存在。

2. 暴露更多並行度

調整Grid/Block配置，目標佔用率最大化：

256線程/塊達到Fermi架構下最佳佔用率（8塊×4 warps=32 warps，接近48上限）。

總結

1. 優先保證對齊與合併訪問，是帶寬利用率的基石
2. 根據訪問模式選擇L1/只讀緩存/繞過L1
3. 採用SoA佈局天然滿足合併訪問
4. 循環展開提升內存操作並行度，效果顯著
5. Grid/Block配置需匹配硬件資源實現滿佔用

參考書：《Professional CUDA C Programming》