Frontier、Privatization 與其他 BFS 最佳化 (Frontiers, Privatization, Other Optimizations)

重點總覽 (Overview)

前兩篇的 vertex-/edge-centric BFS 每個 level 都會對「全部 vertex 或 edge」啟動 thread，浪費大量無用工作。本篇用 frontier (邊界佇列) 只啟動處理「上一層真正活躍 vertex」的 thread，再用 privatization 把 frontier 搬到 shared memory 降低 atomic 競爭，最後處理 launch overhead 與 load balance。

Frontier (Improving Efficiency with Frontiers)

動機

• 無 frontier 版本：每個 level 都對所有 vertex (push) 或所有 edge (edge-centric) 啟動 thread，多數 thread 只發現「我不相關」就結束 → 大量 wasted work。
• Frontier 版本：上一層的 thread 協作建立 frontier，記錄它們新訪問到的 vertex；下一層只對 frontier 內的 vertex 啟動 thread (Luo et al., 2010)。

無 frontier (push)           有 frontier (push)
所有 9 個 vertex 各 1 thread   只有 prevFrontier 內的 vertex
┌─┬─┬─┬─┬─┬─┬─┬─┬─┐          prevFrontier = [1, 2]
│0│1│2│3│4│5│6│7│8│          ┌─┬─┐
└─┴─┴─┴─┴─┴─┴─┴─┴─┘          │1│2│  ← 只啟動 2 個 thread
 只有 1,2 真正做事            └─┴─┘
 其餘 7 個白跑                 currFrontier ← 新訪問的鄰居

Kernel 程式碼 (Fig. 15.12)

__global__ void bfs_frontier_kernel(
        CSRGraph csrGraph, unsigned int* level,
        unsigned int* prevFrontier, unsigned int* currFrontier,
        unsigned int numPrevFrontier, unsigned int* numCurrFrontier,
        unsigned int currLevel) {
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < numPrevFrontier) {                          // 只 launch 上層 frontier 大小
        unsigned int vertex = prevFrontier[i];          // 取出要處理的 vertex
        for (unsigned int edge = csrGraph.srcPtrs[vertex];
             edge < csrGraph.srcPtrs[vertex + 1]; ++edge) {
            unsigned int neighbor = csrGraph.dst[edge];
            // 一次原子完成「檢查未訪問 + 標記」
            if (atomicCAS(&level[neighbor], UINT_MAX, currLevel) == UINT_MAX) {
                unsigned int idx = atomicAdd(numCurrFrontier, 1);  // 取得入隊位置
                currFrontier[idx] = neighbor;                       // 寫入 currFrontier
            }
        }
    }
}

為什麼必須用 atomicCAS (而非 idempotent 寫入)

• atomicCAS(addr, compare, val)：若 *addr == compare 則 *addr = val；回傳舊值。
• 判斷成功：atomicCAS(...) == UINT_MAX 代表「這個 thread 是第一個訪問者」，只有它能入隊。

Privatization (Reducing Contention with Privatization)

概念

與 09-Parallel-Histogram/02-Histogram-Optimizations-Privatization-Coarsening-Aggregation 中的 privatized histogram 相同思路：

• 每個 block 在 shared memory 維護自己的 local frontier 與 local counter。
• thread 只與同 block 的 thread 競爭該 local counter → 競爭範圍從「整個 grid」縮小到「一個 block」。
• 結束時整批 commit 到 global frontier，且寫入可 coalesced (memory coalescing)。

       Block 0                 Block 1                Block 2
  ┌──────────────┐       ┌──────────────┐      ┌──────────────┐
  │ shared mem   │       │ shared mem   │      │ shared mem   │
  │ frontier_s[] │       │ frontier_s[] │      │ frontier_s[] │
  │ count_s (低延遲)│      │ count_s      │      │ count_s      │
  └──────┬───────┘       └──────┬───────┘      └──────┬───────┘
         │  block 內各 thread     │                    │
         │  atomicAdd 只撞同 block │                    │
         ▼                       ▼                    ▼
   一次 atomicAdd(global, count_s) 取得 startIdx，整批 coalesced 寫入
  ┌───────────────────────────────────────────────────────────┐
  │             global currFrontier[]  (public)                │
  └───────────────────────────────────────────────────────────┘

Kernel 程式碼 (Fig. 15.14)

__global__ void bfs_privatized_kernel(/* ...同上參數... */
        unsigned int currLevel) {
    // (1) 宣告並初始化 block 私有 frontier
    __shared__ unsigned int currFrontier_s[LOCAL_FRONTIER_CAPACITY];
    __shared__ unsigned int numCurrFrontier_s;
    if (threadIdx.x == 0) numCurrFrontier_s = 0;
    __syncthreads();                                   // 等 counter 初始化完成

    // (2) BFS 主體 (同 frontier 版)
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < numPrevFrontier) {
        unsigned int vertex = prevFrontier[i];
        for (unsigned int edge = csrGraph.srcPtrs[vertex];
             edge < csrGraph.srcPtrs[vertex + 1]; ++edge) {
            unsigned int neighbor = csrGraph.dst[edge];
            if (atomicCAS(&level[neighbor], UINT_MAX, currLevel) == UINT_MAX) {
                unsigned int idx_s = atomicAdd(&numCurrFrontier_s, 1); // shared，低延遲
                if (idx_s < LOCAL_FRONTIER_CAPACITY) {
                    currFrontier_s[idx_s] = neighbor;                  // 進 local frontier
                } else {                                               // local 溢位 → fallback
                    numCurrFrontier_s = LOCAL_FRONTIER_CAPACITY;       // 復原 counter
                    unsigned int idx = atomicAdd(numCurrFrontier, 1);  // 直接寫 global
                    currFrontier[idx] = neighbor;
                }
            }
        }
    }
    __syncthreads();                                   // 確保 local frontier 已完成

    // (3) 由一個 thread 在 global frontier 一次配置整段空間
    __shared__ unsigned int currFrontierStartIdx;
    if (threadIdx.x == 0)
        currFrontierStartIdx = atomicAdd(numCurrFrontier, numCurrFrontier_s);
    __syncthreads();

    // (4) coalesced 地把 local frontier 搬到 global frontier
    for (unsigned int idx_s = threadIdx.x; idx_s < numCurrFrontier_s;
         idx_s += blockDim.x) {
        currFrontier[currFrontierStartIdx + idx_s] = currFrontier_s[idx_s];
    }
}

Reducing Launch Overhead (Other Optimizations)

• BFS 前幾個 level 的 frontier 通常很小 (第 1 層只有 source 的鄰居)；末幾層也可能很小。
• 對這些小 frontier 的 level，終止舊 grid + 啟動新 grid 的開銷可能超過平行收益。
• 解法:準備一個只用單一 thread block 的 kernel，可連續處理多個 level，level 之間以 __syncthreads() 同步,全程只用 block 級 local frontier。

(A) 每個 level 各 launch 一次 grid     (B) 小 level 合併進單一 block grid
 host ─ launch ─► [grid L0]            host ─ launch ─► ┌ single-block grid ┐
        launch ─► [grid L1]                              │ L0 → __syncthreads()│
        launch ─► [grid L2] ...                          │ L1 → __syncthreads()│
   ▲ 每次都有 launch overhead                            │ (frontier 溢位即返回)│
                                                         └─────────────────────┘
                                       ▲ 一次 launch 跑完多個小 level

• 當 frontier 大到溢出 block-level frontier 時，threads 把 block-level frontier 複製到 global frontier 後返回 host；host 再改用一般 (multi-grid) kernel 繼續，直到 frontier 又變小。
• 效果：消除小 frontier 迭代的 launch overhead (書中留作習題實作)。

Improving Load Balance (Other Optimizations)

• vertex-centric 中每個 thread 的工作量 = 其 vertex 的 degree。社群網路的 celebrity vertex degree 可能比一般 vertex 高出數個數量級 → 少數 thread 拖垮整個 grid。
• 已知對策之一：改用 edge-centric 實作 (見 15-Graph-Traversal/02-Vertex-Centric-and-Edge-Centric-BFS)。
• 另一對策 (Merrill & Garland, 2012)：把 frontier 內 vertex 依 degree 分桶，每桶用合適規模的處理器群:

frontier 內 vertex 依 degree 分 3 桶
  small  degree ──► kernel A: 1 vertex / 1 thread
  medium degree ──► kernel B: 1 vertex / 1 warp   (32 threads)
  large  degree ──► kernel C: 1 vertex / 1 block

Further Challenges & Summary (15.7–15.8)

• BFS 雖是最簡單的圖演算法，卻體現複雜應用的共同挑戰：
  1. problem decomposition 萃取平行性 (vertex vs edge vs frontier)。
  2. 善用 privatization 降低 atomic 競爭。
  3. 實作 fine-grained load balancing (分桶)。
  4. 確保正確 synchronization (__syncthreads()、atomicCAS)。
• 進階挑戰：圖大小超過 GPU 記憶體容量；以及前處理成其他格式以暴露更多平行性 / locality / 利於 load balance。
• 應用領域:推薦系統、社群偵測、圖樣式比對、異常偵測。

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

Related Notes

• 15-Graph-Traversal/01-Graph-Representations-and-BFS
• 15-Graph-Traversal/02-Vertex-Centric-and-Edge-Centric-BFS
• 09-Parallel-Histogram/02-Histogram-Optimizations-Privatization-Coarsening-Aggregation
• 09-Parallel-Histogram/01-Atomic-Operations-and-Basic-Histogram
• 06-Performance-Considerations/01-Memory-Coalescing
• 06-Performance-Considerations/03-Thread-Coarsening