CUDA 並行求和性能分析：共享內存歸約優化與內存傳輸瓶頸

實驗背景

測試數組大小 2²⁰ (1,048,576) 元素，重複 1,000 次求和，對比 CPU、GPU（含/不含內存分配拷貝）性能。

測試結果

單次運行微觀對比

• CPU：~10 ms
• GPU 含分配/拷貝：~60 ms（慢於 CPU）
• GPU 僅核函數：~0 ms（微秒級）

核心結論

1. GPU 計算極快：歸約核函數利用共享內存並行樹狀求和，計算吞吐遠超 CPU。
2. 內存分配與拷貝是瓶頸：cudaMalloc、cudaMemcpy 開銷巨大，特別是大數據量時。
3. 設計原則：
   • 數據盡量常駐 GPU 內存，避免反覆 Host↔Device 拷貝。
   • 批次處理：一次拷貝，多次核函數調用。
   • 若數據源在 CPU 且僅用一次，GPU 可能不划算。

優化建議

• 使用 pinned memory (cudaMallocHost) 加速拷貝。
• 異步拷貝 cudaMemcpyAsync + 流 cudaStream 實現計算與傳輸重疊。
• 持久化內核 / Graph Capture 減少啟動開銷。
• 大規模歸約考慮多級歸約（如本文後續部分的迭代優化）。

代碼關鍵點回顧

// 樹狀歸約核函數（共享內存）
__global__ void block_sum(const double* input, double* per_block_results, size_t n) {
    extern __shared__ double sdata[];
    size_t i = blockIdx.x * blockDim.x + threadIdx.x;
    double x = (i < n) ? input[i] : 0;
    sdata[threadIdx.x] = x;
    __syncthreads();
    for (int offset = blockDim.x/2; offset > 0; offset >>= 1) {
        if (threadIdx.x < offset) sdata[threadIdx.x] += sdata[threadIdx.x + offset];
        __syncthreads();
    }
    if (threadIdx.x == 0) per_block_results[blockIdx.x] = sdata[0];
}

總結

GPU 並行歸約計算本身極其高效，但端到端性能受限於內存傳輸。實際工程中必須從數據佈局、生命週期、傳輸策略全局優化，而非單純優化核函數。對於一次性小數據任務，CPU 往往更優；對於反覆計算、大數據、數據已在 GPU 的場景，GPU 優勢壓倒性。