前言
線程組織形式直接決定 GPU 並行效率。本文以矩陣加法為例，詳細演示 2D Grid + 2D Block 映射 2D 矩陣的索引計算、記憶體存取模式，並透過實測對比不同 Block 尺寸對效能的影響。
矩陣線性儲存與線程映射
矩陣按行優先儲存：A[row * N + col]。
2D 索引計算
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
int idx = row * N + col; // 全域線性索引
邊界檢查：Kernel 入口必須驗證 row < N && col < N，防止越界存取。
完整 Kernel 實現
__global__ void sumMatrix2D(const float *A, const float *B, float *C, int N) {
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
if (row < N && col < N) {
int idx = row * N + col;
C[idx] = A[idx] + B[idx];
}
}
Host 啟動配置
int N = 1 << 14; // 16384 x 16384
dim3 block(32, 32); // 1024 線程/Block
dim3 grid((N + block.x - 1) / block.x, (N + block.y - 1) / block.y);
sumMatrix2D<<>>(d_A, d_B, d_C, N);
實測效能對比
| Block 尺寸 | Grid 尺寸 | 耗時 (秒) | 分析 |
|---|
| 32×32 (1024) | 512×512 | 0.0603 | 基準 |
| 32×16 (512) | 512×1024 | 0.0380 | Block 數翻倍，SM 佔用率提升，效能最佳 |
| 16×16 (256) | 1024×1024 | 0.0455 | Block 過多導致調度開銷上升，效能下降 |
效能分析核心因素
- 佔用率：32×16 產生更多 Block，更好填滿 SM。
- Warp 效率：所有配置均為 32 倍數，無非滿 Warp 損耗。
- 記憶體合併：連續 32 線程存取連續 32 floats → 完美 128-byte 合併事務。
- 調度開銷：過多 Block (1024×1024) 增加調度器負擔。
最佳實踐清單
- ✅ Block 尺寸為 32 倍數（128~1024）。
- ✅ 2D/3D 問題優先用 2D/3D Block，索引計算自然對齊記憶體存取。
- ✅ Grid 尺寸 ≥ 2×SM 數，確保負載均衡。
- ✅ 實測調優：
nvprof --metrics achieved_occupancy,elapsed_cycles。
- ⚠️ 避免「越大越好」，存在最優平衡點。
擴展：非方陣與批次處理
