執行緒粗化 (Thread Coarsening)

重點總覽 (Overview)

細粒度 vs 粗化 (Finest Granularity vs Coarsening)

到目前為止的所有 kernel 都採用最細粒度平行化（每 thread 最小工作單位）：

最細粒度的優點 — transparent scalability

• 暴露最多平行度給硬體。
• 資源夠 → 全部平行執行；資源不夠 → 硬體自動把 thread blocks 一個接一個串行化，程式不必改。

最細粒度的缺點 — price of parallelism

• 平行化常需付出額外成本：
  • redundant loading of data（不同 block 重複載入同一份資料）
  • redundant work（重複計算）
  • synchronization overhead（同步開銷）
• thread 真的平行跑時，這個 price 值得付。
• 但若硬體因資源不足而串行化了這些工作 → price 白付了。

範例：Tiled 矩陣乘法 (Coarsening for Tiled MatMul)

Price of parallelism = 重複載入 M tile

計算 P 中兩個水平相鄰的 output tile 時：

                N (input, 不同 tile)
            ┌──────────┬──────────┐
            │  N tile  │  N tile  │   ← 每個 output tile 要不同的 N tile
            │   (左)   │   (右)   │
   ─────────┼──────────┼──────────┤
   M tile   │   P_tile │  P_tile  │   ← 兩個相鄰的 output tile
  (同一塊!) │    左    │    右    │
   ─────────┴──────────┴──────────┘
       ▲
       └─ 同一塊 M input tile 同時餵給「左」「右」兩個 output tile

• 兩個 output tile 需要不同的 N input tiles。
• 但需要相同的 M input tiles。
• Ch.5 tiled 版：每個 output tile 由不同 block 處理 → shared memory 不能跨 block 共享 → 每個 block 各載入一份 M tile（重複！）。

Coarsening loop 程式碼 (Fig. 6.13)

#define TILE_WIDTH   32
#define COARSE_FACTOR 4          // line 02: 每個 coarsened thread 負責幾個 work unit

__global__ void matrixMulKernel(float* M, float* N, float* P, int Width) {
    __shared__ float Mds[TILE_WIDTH][TILE_WIDTH];
    __shared__ float Nds[TILE_WIDTH][TILE_WIDTH];

    int bx = blockIdx.x;  int by = blockIdx.y;
    int tx = threadIdx.x; int ty = threadIdx.y;

    int row = by * TILE_WIDTH + ty;
    // line 13: 第一個負責的 column；block 跨度 = TILE_WIDTH * COARSE_FACTOR
    int colStart = bx * TILE_WIDTH * COARSE_FACTOR + tx;

    // lines 16-19: 為每個負責的元素各開一個 Pvalue 累加器
    float Pvalue[COARSE_FACTOR];
    for (int c = 0; c < COARSE_FACTOR; ++c)      // line 17: 初始化的 coarsening loop
        Pvalue[c] = 0.0f;

    for (int ph = 0; ph < Width / TILE_WIDTH; ++ph) {   // line 22: 走訪 input tiles
        // 一個 phase 只載入「一塊」M tile（與原版相同）
        Mds[ty][tx] = M[row * Width + ph * TILE_WIDTH + tx];

        for (int c = 0; c < COARSE_FACTOR; ++c) {       // line 27: 主 coarsening loop
            int col = colStart + c * TILE_WIDTH;        // line 29: 算出本次負責的 column

            Nds[ty][tx] = N[(ph * TILE_WIDTH + ty) * Width + col];  // line 32: 載入對應 N tile
            __syncthreads();

            for (int k = 0; k < TILE_WIDTH; ++k)        // lines 35-37: 用同一塊 M tile
                Pvalue[c] += Mds[ty][k] * Nds[k][tx];   //   更新「不同」的 Pvalue[c]
            __syncthreads();
        }
    }

    for (int c = 0; c < COARSE_FACTOR; ++c) {           // lines 44-47: 寫回各輸出元素
        int col = colStart + c * TILE_WIDTH;
        P[row * Width + col] = Pvalue[c];
    }
}

關鍵改動對照

算術強度提升 (Arithmetic Intensity, OP/B)

以 floating-point 對 global memory 存取比 (OP/B) 量化好處（書末習題 4）：

• Coarsening 把每組 4 個 output tile 的 M tile 載入從 4 塊降到 1 塊 → global memory traffic 下降 → OP/B 從 8 升到 12.8。

三大陷阱 (Three Pitfalls)

跨章節的 coarsening 應用 (Where It Recurs)

不同章節「price of parallelism」不同，coarsening 攻擊點也不同：

考試/面試重點 (Exam / Test Patterns)

Related Notes

• 05-Memory-Architecture-And-Data-Locality/02-Tiling-and-Tiled-Matrix-Multiplication
• 06-Performance-Considerations/01-Memory-Coalescing
• 06-Performance-Considerations/02-Hiding-Memory-Latency
• 06-Performance-Considerations/04-Optimization-Checklist-and-Bottlenecks
• 08-Stencil/03-Thread-Coarsening-and-Register-Tiling
• 10-Reduction/03-Scaling-Reduction-Hierarchical-and-Coarsening