Tiled 與 Circular Buffer Merge Kernel (Memory Coalescing)

重點總覽 (Overview)

本筆記延續 12-Merge/02-Co-Rank-Function-and-Basic-Merge-Kernel，解決 basic merge kernel 的記憶體存取弱點：透過 shared memory tiling 取得 coalesced load，再用 circular buffer 消除 tiled 版本的頻寬浪費。

Basic Kernel 的兩個 Coalescing 弱點 (Motivation)

basic merge kernel 每個 thread 獨立做 sequential merge，導致：

warp 中相鄰三個 thread (t0,t1,t2) 在 merge_sequential 第一輪:
  讀取 →  A[0]      A[2]      B[0]      ← 位址不連續，未 coalesced
  寫入 →  C[0]      C[3]      C[6]      ← 位址不連續，未 coalesced

1. Sequential merge 階段：相鄰 thread 的 input/output 子陣列彼此相鄰，但每個 thread 內部是順序掃描，同一時刻的存取位址分散 → 未 coalesced。
2. Co-rank 階段：binary search 本身存取模式不規則，且直接打 global memory → 幾乎不可能 coalesced。

Block-Level Co-Rank 與協同載入 (Block-Level Co-Rank & Coalesced Loading)

關鍵觀察：一個 block 內所有 thread 共同使用 A、B 的較大「block 級子陣列」生成一個較大的 C 子陣列。因此只需整個 block 呼叫一次 co-rank 取得 block 級邊界，全 block 即可協同 coalesced 載入。

        block 0          block 1          block 2
C: |■■■■■■■■|--------|■■■■■■■■|--------|■■■■■■■■|   ← C 平均切成 gridDim.x 段
        ↑ co_rank(C_curr) / co_rank(C_next) 決定 A、B 的對應切點
A: |■■■■■■■■■■■■|----|■■■|--------------|■■■■■■|   ← 各 block 的 A 子陣列大小隨資料而異
B: |■■|--------------|■■■■■■■■■■■■|------|■■■■|     ← 但 A_len + B_len = C_len 恆成立

Part 1：只用 1 個 thread 算 block 級 co-rank，存入 shared memory 給全 block 看（減少 global 存取）：

__global__ void merge_tiled_kernel(int* A, int m, int* B, int n,
                                   int* C, int tile_size) {
    extern __shared__ int shareAB[];
    int* A_S = &shareAB[0];           // 前半: tile_size 個 A 元素
    int* B_S = &shareAB[tile_size];   // 後半: tile_size 個 B 元素
    int C_curr = blockIdx.x * ceil((m+n)/(float)gridDim.x);
    int C_next = min((blockIdx.x+1) * (int)ceil((m+n)/(float)gridDim.x), m+n);
    if (threadIdx.x == 0) {                       // 只 1 個 thread 做 co-rank
        A_S[0] = co_rank(C_curr, A, m, B, n);     // 借 shared memory 廣播給全 block
        A_S[1] = co_rank(C_next, A, m, B, n);
    }
    __syncthreads();                              // 等 co-rank 結果可見
    int A_curr = A_S[0], A_next = A_S[1];
    int B_curr = C_curr - A_curr, B_next = C_next - A_next;
    __syncthreads();
    // ... 接 Part 2
}

迭代式 Tile 載入與 Consumed 追蹤 (Iterative Loading with Consumed Tracking)

每個 block 需 total_iteration = ceil(C_length / tile_size) 輪。每輪協同載入 x 個 A + x 個 B（x = tile_size），用 threadIdx.x 選元素使存取 coalesced。

Part 2（載入 tile）：

    int counter = 0;
    int C_length = C_next - C_curr;
    int A_length = A_next - A_curr;
    int B_length = B_next - B_curr;
    int total_iteration = ceil(C_length / (float)tile_size);
    int C_completed = 0, A_consumed = 0, B_consumed = 0;
    while (counter < total_iteration) {
        // blockDim.x 個 thread 每輪載 blockDim.x 個 → 共跑 tile_size/blockDim.x 次
        for (int i = 0; i < tile_size; i += blockDim.x)
            if (i + threadIdx.x < A_length - A_consumed)         // 防止越界 (最後一輪)
                A_S[i+threadIdx.x] = A[A_curr + A_consumed + i + threadIdx.x];
        for (int i = 0; i < tile_size; i += blockDim.x)
            if (i + threadIdx.x < B_length - B_consumed)
                B_S[i+threadIdx.x] = B[B_curr + B_consumed + i + threadIdx.x];
        __syncthreads();
        // ... 接 Part 3

Part 3（thread 級 merge + 更新 consumed）：

        int c_curr = threadIdx.x * (tile_size/blockDim.x);
        int c_next = (threadIdx.x+1) * (tile_size/blockDim.x);
        c_curr = min(c_curr, C_length - C_completed);     // 夾住最後一輪
        c_next = min(c_next, C_length - C_completed);
        int a_curr = co_rank(c_curr, A_S, min(tile_size, A_length-A_consumed),
                                     B_S, min(tile_size, B_length-B_consumed));
        int b_curr = c_curr - a_curr;
        int a_next = co_rank(c_next, A_S, min(tile_size, A_length-A_consumed),
                                     B_S, min(tile_size, B_length-B_consumed));
        int b_next = c_next - a_next;
        merge_sequential(A_S+a_curr, a_next-a_curr, B_S+b_curr, b_next-b_curr,
                         C + C_curr + C_completed + c_curr);   // co-rank 全在 shared memory
        counter++;
        C_completed += tile_size;
        A_consumed += co_rank(tile_size, A_S, tile_size, B_S, tile_size);
        B_consumed = C_completed - A_consumed;   // 由 C_completed 與 A_consumed 導出
        __syncthreads();
    }
}

A_consumed 指標如何推進（block 0 的第 1 輪）：

A 子陣列:  [□□□|■■■■■■■■■■■■......]
            ↑   ↑
   iter0 消耗  A_curr + A_consumed  ← iter1 從這裡開始載入
   (白色=已消耗, 長度=A_consumed)
下一輪 tile 起點 = A[A_curr + A_consumed] , B[B_curr + B_consumed]

Tiled Kernel 的 50% 頻寬浪費 (The Half-Bandwidth Deficiency)

每輪載入:  A_S [x 個] + B_S [x 個]  = 2x 載入
每輪消耗:  只用掉合計 x 個
          → 剩下的 x 個被「丟棄並在下一輪重新載入」
利用率 =  x / 2x = 50%  ← 浪費一半 memory bandwidth

Circular Buffer Merge Kernel (Circular Buffer Scheme)

核心改動：不再每輪從 A_S[0] 重載整個 tile，而是保留上一輪未消耗的元素，只 refill 被消耗掉的部分。新增 A_S_start / B_S_start（tile 在 buffer 中的動態起點）與 A_S_consumed / B_S_consumed（本輪消耗量、即下輪需補的量）。

(A) iter0 載入: A_S_start=0    [■■■■■■■■]              整段新載入
(B) iter0 結束: A_S_start 前移 [□□□|■■■■■]              白=已消耗, 灰=保留
(C) iter1 補載:               [▓▓▓|■■■■■]              只補 A_S_consumed 個 (wrap around)
(D) iter1 結束: A_S_start wrap [■■|□□□|▓▓]              消耗段繞回開頭, start 用 % 環繞
       index 環繞:  idx % tile_size

Part 1 與 tiled 完全相同。Part 2 的差異在迴圈出口與 wrap-around index：

    int A_S_start = 0, B_S_start = 0;
    int A_S_consumed = tile_size, B_S_consumed = tile_size; // 第一輪需填滿
    while (counter < total_iteration) {
        for (int i = 0; i < A_S_consumed; i += blockDim.x)        // 只補需要的量
            if (i+threadIdx.x < A_length-A_consumed && i+threadIdx.x < A_S_consumed)
                A_S[(A_S_start + (tile_size - A_S_consumed) + i + threadIdx.x) % tile_size]
                    = A[A_curr + A_consumed + i + threadIdx.x];   // % 環繞回開頭
        // ... B_S 同理 (留作習題)

Part 3 末尾的更新（環繞 start、更新 consumed）：

        counter++;
        A_S_consumed = co_rank_circular(min(tile_size, C_length-C_completed),
                          A_S, min(tile_size, A_length-A_consumed),
                          B_S, min(tile_size, B_length-B_consumed),
                          A_S_start, B_S_start, tile_size);
        B_S_consumed = min(tile_size, C_length-C_completed) - A_S_consumed;
        A_consumed += A_S_consumed;
        C_completed += min(tile_size, C_length-C_completed);
        B_consumed  = C_completed - A_consumed;
        A_S_start = (A_S_start + A_S_consumed) % tile_size;   // 環繞推進起點
        B_S_start = (B_S_start + B_S_consumed) % tile_size;
        __syncthreads();
    }

Simplified Model 與 Circular 函式 (Simplified Co-Rank Model)

直接把 circular buffer 的真實索引交給 user code 會很複雜（例如 a_next 可能因環繞而小於 a_curr，長度要分兩種情況算 a_next-a_curr 或 a_next-a_curr+tile_size）。

(A) 真實 circular buffer:  b_next 環繞 → 在實體上 b_next < b_curr
    B_S: [ ...b_next | (空) | B_S_start...b_curr... ]
(B) 簡化模型 (user 看到的): 一段連續 tile_size，b_next ≥ b_curr
    B_S: [ B_S_start ............ b_curr .. b_next .. ]

co_rank_circular：邏輯與原 co-rank 完全相同，唯一差別是 i, i-1, j, j-1 不再直接當索引，而是當 offset 加到 start 上、再判斷是否需環繞：

__device__ int co_rank_circular(int k, int* A, int m, int* B, int n,
                                int A_S_start, int B_S_start, int tile_size) {
    int i = (k < m) ? k : m;
    int j = k - i;
    int i_low = max(0, k-n), j_low = max(0, k-m);
    bool active = true;
    while (active) {
        int i_cir     = (A_S_start + i     >= tile_size) ? A_S_start+i-tile_size     : A_S_start+i;
        int i_m_1_cir = (A_S_start + i - 1 >= tile_size) ? A_S_start+i-1-tile_size   : A_S_start+i-1;
        int j_cir     = (B_S_start + j     >= tile_size) ? B_S_start+j-tile_size     : B_S_start+j;
        int j_m_1_cir = (B_S_start + j - 1 >= tile_size) ? B_S_start+j-1-tile_size   : B_S_start+j-1;
        if (i > 0 && j < n && A[i_m_1_cir] > B[j_cir]) {        // i 太大
            int delta = (i - i_low + 1) >> 1;
            j_low = j; j += delta; i -= delta;
        } else if (j > 0 && i < m && B[j_m_1_cir] >= A[i_cir]) {// j 太大
            int delta = (j - j_low + 1) >> 1;
            i_low = i; i += delta; j -= delta;
        } else active = false;
    }
    return i;
}

merge_sequential_circular 同理：邏輯與 Fig.12.2 的順序 merge 一致，只在四處存取 A/B 時把 i、j 轉成 i_cir、j_cir 並判斷環繞。

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

Related Notes

• 12-Merge/02-Co-Rank-Function-and-Basic-Merge-Kernel
• 12-Merge/04-Thread-Coarsening-and-Summary
• 12-Merge/01-Merge-Foundations-and-Co-Rank-Concept
• 05-Memory-Architecture-And-Data-Locality/02-Tiling-and-Tiled-Matrix-Multiplication
• 06-Performance-Considerations/01-Memory-Coalescing
• 07-Convolution/03-Tiled-Convolution-and-Halo-Handling