有序合併與 Co-rank 概念 練習題 (Practice - Merge Foundations and the Co-rank Concept)

Part 1:while-loop 用雙指標 i、j 同時掃描 A、B,每次迭代比較後填一個 C 位置(A[i]<=B[j] 取 A,否則取 B)。Part 2:當一方耗盡後,複製另一方剩餘元素到 C。每個輸入元素讀一次、每個 C 位置寫一次 → O(m+n),與總元素數線性。

用 <= 表示「平手時 A 元素優先放入 C」,維持 stability across lists(跨清單穩定性)——可保留先前依其他 key 排序的成果。若改成 <,平手時改取 B,變成「B 優先」:仍是合法排序,但不再 stable across lists。

rank k = 輸出元素 C[k] 的索引;co-rank (i, j) = 產生 prefix C[0..k-1] 所需的 A、B prefix 長度,即 C[0..k-1] = merge(A[0..i-1], B[0..j-1])。不變式:k = i + j,且 Siebert & Träff 證明對每個 k 此 (i,j) 唯一。co_rank 只回傳 i,caller 自行算 j = k - i。

i = min(k, m)(i 不可超過 A 大小或 k);i_low = max(0, k-n);j_low = max(0, k-m);j = k - i。i_low 設成 0 仍正確但較慢——因 B 最多貢獻 n 個元素,當 k > n 時 A 至少要貢獻 k-n 個,故下界可收緊為 k-n 以縮小搜尋範圍。

接受條件:A[i-1] <= B[j] 且 B[j-1] < A[i](注意 <=/< 不對稱以維持 stability)。delta = (i - i_low + 1) >> 1 是 ceil,確保每步至少縮減 1,避免搜尋區間卡住、保證 while-loop 終止。搜尋為 O(log₂N)。

第一次以 k_curr 為 rank → 得 i_curr(輸入子陣列起點);第二次以 k_next 為 rank → 得 i_next(輸入子陣列終點)。A 子陣列長度 = i_next - i_curr,B 長度 = j_next - j_curr。k_next 須 min((tid+1)*elems, m+n) 做 clamp,因 m+n 不一定整除 thread 數。

直接把環繞索引交給 user code 很複雜(a_next 可能因 wrap 而小於 a_curr,長度要分兩種情況算)。Simplified model 對外呈現「tile 從 A_S_start 起、連續最多 tile_size 個」的幻覺,使 a_next ≥ a_curr 恆成立;環繞只在實際存取元素時於 _circular 函式內展開。co_rank_circular 多 3 個參數:A_S_start, B_S_start, tile_size。

因為輸入用量是 data-dependent:最壞情況下,當前 tile_size 個輸出可能全部來自單一輸入陣列(全 A 或全 B)。為保證每個 thread 都拿得到所需元素,A、B 必須各備 tile_size 個。實際只會消耗合計 tile_size 個(A_consumed_iter + B_consumed_iter = tile_size),剩下一半被丟棄重載 → tiled kernel 浪費 50% 頻寬。

合併結果 C=[1,7,7,8,9,10,10,10,12]。C[8] 的 rank k=8,prefix C[0..7]=[1,7,7,8,9,10,10,10] = merge(A[0..4], B[0..2]),故 co-rank i=5(整個 A)、j=3,驗證 i+j=5+3=8 ✓(C[8]=12 來自 B[3])。

總輸出 m+n = 1,638,400,thread 數 = 1,638,400 / 8 = 204,800。
(a) basic:每個 thread 都在 global memory 做 → 204,800。
(b) tiled:每 block 只 1 個 thread 做 block 級 global co-rank → blocks = 204,800/1024 = 200 → 200。
(c) tiled:thread 級 co-rank 全在 shared memory → 全部 204,800。

Shared memory = A_S 與 B_S 各 1024 個 int → (1024+1024)×4 = 8192 bytes。每 block 輸出 = 64000/16 = 4000,迭代次數 = ceil(C_length/tile_size) = ceil(4000/1024) = 4 輪。每輪載入 for-loop 跑 tile_size/blockDim.x = 1024/128 = 8 次。

Basic:A/B/C 全在 global,每 thread 順序掃描造成相鄰 thread 讀 A[0],A[2],B[0]、寫 C[0],C[3],C[6] → uncoalesced,co-rank binary search 也不規則。Tiled:用 Ch.6 策略(3) coalesced 載入 shared memory,但每輪載 2×tile_size 只用一半 → 50% 利用率。Circular buffer:保留上輪未消耗元素只 refill 被消耗部分(% tile_size 環繞)→ ~100% 利用率。由於 merge 是 memory-bandwidth bound,節省頻寬直接提升效能,值得付出程式複雜度與額外 register 的代價。

Merge 平行化的主要代價是每個 thread 都要做自己的 binary-search co-rank(輸入範圍 data-dependent)。T=1(未 coarsen)時每個輸出元素都要一次 binary search → prohibitively expensive;coarsening 讓 (m+n)/T 個 thread 每個只做 2 次 co-rank,中間 T 個元素用 O(T) 的 merge_sequential 填滿,把 2·log N 的 search 攤提到 T 個元素。不同點:reduction/scan 的 coarsening 省的是同步 / 重複載入 / atomic 成本,merge 省的是 binary search 次數。代價:T 太大 → thread/block 數變少 → occupancy / 平行度下降(但 memory-bound 下通常可接受,最佳 T 需 profiling)。

核心觀念鏈:輸入範圍 data-dependent(dynamic input identification)→ co_rank binary search O(log N) 反查 → thread coarsening 攤提 search 成本 → shared-memory tiling 取得 coalesced load → circular buffer 達 100% 頻寬利用 → simplified model 把環繞複雜度封裝進 library 函式。

關鍵公式:k = i + j(co-rank 不變式)、i = min(k,m)、i_low = max(0,k-n)、j_low = max(0,k-m)、delta = (i-i_low+1)>>1、total_iteration = ceil(C_length/tile_size)、numThreads = (m+n)/T。

一句話總結:merge 是 memory-bandwidth bound 的「動態輸入辨識」pattern,co-rank 解決「要讀哪段」,tiling + circular buffer 解決「怎麼讀得快」,coarsening 解決「search 成本太貴」。