Enter your search text in the box above

Select a result to preview

問題分解：Output-centric 與 Input-centric (Problem Decomposition)

#computational-thinking #scatter-vs-gather #load-balancing #atomic-operations #problem-decomposition

重點總覽 (Overview)

問題分解 (problem decomposition) 是平行程式三步驟（algorithm selection → problem decomposition → optimization/tuning）的第二步：把問題拆成可同時安全求解的子問題，關鍵是決定每個平行執行單元要做哪份工作。最常見的兩種策略：

策略 thread 對應到 存取模式 累積位置 atomics? GPU 偏好
Output-centric 一個 output 元素 Gather（收集多個 input） private register 否 通常較佳
Input-centric 一個 input 元素 Scatter（散佈到多個 output） global memory（眾人可寫） 是 通常較差
Important

兩種分解結果相同，但在特定硬體上效能可能天差地別。選擇不能只看 gather/scatter，還要綜合：平行度 (parallelism)、input→output 映射難易、load balance。

兩種分解策略 (Output-centric vs Input-centric)

(A) OUTPUT-CENTRIC = GATHER          (B) INPUT-CENTRIC = SCATTER
  in0 in1 in2 in3   in0 in1 in2       T1=in0        T2=in1
    \   \  |  /        \  |  /        / | \ \       / | \ \
     \   \ | /          \ | /        v  v  v v     v  v  v v
      Thread1            Thread2    out0 out1 ...  out0 out1 ...
       (out0)            (out1)        ^----同一 out 多人寫----^
        |                  |           需 atomicAdd 防 race condition
      out[0]   <-register acc->        (多個 thread 同時更新同一格)
   單一寫者，免 atomics            慢：atomics << register 存取

Gather access（output-centric 的優勢）

Scatter access（input-centric 的代價）

// Output-centric (gather): 一個 thread 對應一個 OUTPUT
__global__ void gatherKernel(float* out, const float* in) {
    int o = blockIdx.x * blockDim.x + threadIdx.x;     // output 索引
    if (o < N_OUT) {
        float acc = 0.0f;                              // 累積在 register
        for (int k = 0; k < numInputsFor(o); ++k)      // 掃過貢獻的 inputs
            acc += contribution(in[inputOf(o, k)], o);
        out[o] = acc;                                  // 單一寫者 -> 免 atomics
    }
}

// Input-centric (scatter): 一個 thread 對應一個 INPUT
__global__ void scatterKernel(float* out, const float* in) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;     // input 索引
    if (i < N_IN) {
        for (int k = 0; k < numOutputsFor(i); ++k)     // 掃過受影響的 outputs
            atomicAdd(&out[outputOf(i, k)],            // 多人寫 -> 必須 atomics
                      contribution(in[i], k));
    }
}

選擇分解策略的考量 (Decision Criteria)

除了 gather vs scatter（atomics），還要看：

考量 說明 偏好
Gather vs Scatter scatter 需慢速 atomics 多偏 gather (output-centric)
平行度 (parallelism exposed) 哪邊的元素數量多、能餵飽更多 thread 元素數較多的那邊
映射難易 (input→output mapping) thread 是否容易知道自己負責哪些對應元素 容易識別者
Load balance 每個 thread 工作量是否平均 工作量均勻者
Warning

「output-centric 永遠較好」是錯誤簡化。當輸出元素太少（平行度不足）、或輸出端工作量嚴重不均（load imbalance）、或 mapping 難以判定時，input-centric 反而較佳（histogram、SpMV/COO 即是）。最佳分解常取決於 input dataset。

Tip

當 update 是 idempotent（重複寫同值無害，如 BFS 的 level 設定）時，scatter 不需要 atomics，此時 gather/scatter 不再是考量點，只剩平行度與 load balance 決定勝負（graph traversal 即是）。

全書平行模式分解策略巡禮 (Survey Across the Book)

計算 / 章節 分解 存取 為何如此選擇
Image / MatMul / Convolution / Stencil (Ch3,5,6,7,8) Output-centric Gather thread 對應 output（pixel/矩陣元素/grid 點）並掃 input；輸出元素夠多（平行度足）、mapping 簡單、各輸出工作量相同（無 load imbalance）→ gather 免 atomics 全勝
Histogram (Ch9) Input-centric Scatter (atomics) output-centric 會：①bin 數 << input 數（平行度大減）②thread 不掃過全部 input 就不知哪些屬於它（不 work-efficient）③各 bin input 數不同（load imbalance）→ 三點皆不利，故選 input-centric
Merge (Ch12) Output-centric （gather 非主因） 每個 output 只由一個 input 貢獻，input-centric 其實也不需 atomics；但 input-centric 要嘛使 input→output 映射昂貴，要嘛造成高 load imbalance，故仍選 output-centric
SpMV/COO (Ch14) Input-centric Scatter (atomics) thread 對應 nonzero、atomically 更新 output vector；雖用 atomics，卻能抽取更多平行度且 load balance 更好
SpMV 其他 kernel (Ch14) Output-centric Gather thread 對應 output vector 元素並掃 nonzeros
Hybrid ELL-COO (Ch14) 混合 mixed output-centric + input-centric 混合分解的範例；最佳取決於資料集
Graph: vertex push / edge-centric (Ch15) Input-centric (idempotent, 免 atomics) thread 對應 vertex/edge，更新鄰居 level；scatter/gather 非考量，由平行度與 load balance 決定，最佳依資料集
Graph: vertex pull (Ch15) Output-centric (idempotent) thread 只更新自己負責的 vertex level
MRI FhD (Ch17) Output-centric Gather 每個 thread 用 gather 算一個 FhD 元素；避免 atomics；每個 k-space sample 被重複使用多次
Electrostatic potential (Ch18) Output-centric (grid-centric) Gather grid-centric（一 thread 算所有 atom 對一 grid 點的貢獻）優於 atom-centric（input-centric）；gather 較佳、平行度足夠
Important

Electrostatic potential 的兩個困難如何被克服（output-centric/grid-centric 的代價）：

  1. input→output 映射難（grid 點不知哪些 atom 在 cutoff 內）→ 用 binning 解決。
  2. binning 造成 load imbalance（每 bin atom 數不同）→ 用 SmallBin-Overlap 實作（把序列式 overflow 處理與下個 kernel 重疊）解決。

關鍵量化 (Key Quantities)

量 Output-centric Input-centric
平行度 (#threads) ≈ #output 元素 ≈ #input 元素
Histogram 工作複雜度 O(#bins × #inputs)（每 thread 須掃全部 input → 不 work-efficient） O(#inputs)（每 input 處理一次 → work-efficient）
同步成本 register 累積，免 atomics atomicAdd（慢）
Warning

Histogram 的 output-centric 之所以不可行，正是因為其工作量為 O(#bins × #inputs) 而非 O(#inputs)：缺少快速「哪些 input 屬於此 bin」的索引，每個 bin thread 都得掃過所有 input。

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

情境 / 關鍵字 答案 / 技巧
"gather access pattern" Output-centric；每 thread 收集多個 input 成一個 output，累積在 register，可用 constant/shared memory 共享 input，免 atomics（GPU 偏好）
"scatter access pattern" Input-centric；每 thread 把一個 input 散佈更新到多個 output，需 atomic operations 防 race condition，較慢
為何 GPU 通常偏好 output-centric 結果存 private register、可共享 input 省 bandwidth、避免慢的 atomics
Histogram 為何用 input-centric output-centric 三大缺點：平行度低（bin 少）、不 work-efficient（O(bins×inputs)）、load imbalance
Merge 為何 output-centric（但 gather 非理由） 每 output 只來自一個 input，input-centric 也免 atomics；但 input-centric 會 mapping 貴 或 load imbalance
SpMV/COO 用 input-centric 的代價/好處 代價=atomics；好處=更高平行度 + 更好 load balance；最佳依資料集
Graph BFS push vs pull push/edge = input-centric，pull = output-centric；level 更新 idempotent → 免 atomics，故看平行度與 load balance
選分解策略要考量哪些面向 ①gather vs scatter (atomics) ②平行度 ③input→output 映射難易 ④load balance
Hybrid ELL-COO 代表什麼 output-centric 與 input-centric 的混合分解
EP grid-centric 兩困難如何解 mapping 難 → binning；load imbalance → SmallBin-Overlap
兩種分解結果是否相同 相同結果，但效能可能差很多；非單一策略恆優