Vertex-centric 與 Edge-centric 的 BFS 平行化 (Push, Pull, Edge-Centric)
重點總覽 (Overview)
每一層 (level) 都呼叫一次 kernel:先把 root 標為 level 0,逐層 (level-by-level) 把上一層鄰居標為當層,直到沒有新頂點被造訪。本筆記比較三種把工作映射到 thread 的策略。
| 策略 | Thread 對應 | 需要的格式 | 檢查條件 | 動作 | 別名 |
|---|---|---|---|---|---|
| Vertex-centric Push | 1 thread / vertex | CSR(出邊 outgoing) | 我的 vertex 在 currLevel-1? |
把所有未造訪鄰居標為 currLevel |
top-down / push |
| Vertex-centric Pull | 1 thread / vertex | CSC(入邊 incoming) | 我的 vertex 未造訪? | 若有鄰居在 currLevel-1,標自己並 break |
bottom-up / pull |
| Edge-centric | 1 thread / edge | COO(src+dst) | src 在 currLevel-1 且 dst 未造訪? |
把 dst 標為 currLevel |
— |
| 比較維度 | Push (CSR) | Pull (CSC) | Edge-centric (COO) |
|---|---|---|---|
| 暴露的平行度 | 中(= V) | 中(= V) | 高(= E,最多) |
| 早期 level 效率 | 好(只跑少數前層頂點) | 差(所有未造訪頂點都跑) | 差(掃全部邊) |
| 後期 level 效率 | 差 | 好(早退 break) | 差 |
| Load imbalance / divergence | 高(依 degree) | 中(可 break) | 最低(每 thread 恰 1 邊) |
| 儲存成本 | 低 | 低 | 高(src+dst 兩陣列) |
| 適用 | 早期層、低 degree | 後期層、高 degree | 小圖 / degree 變異大的圖 |
同步必須跨 level:每層都重新 launch 一個 grid,因為要等上一層全部標完才能標下一層;若在同一 grid 內跨層處理,可能把頂點標到錯誤的 level。
Vertex-centric Push(top-down / CSR)
- 每個 thread 負責一個 vertex,先檢查自己是否屬於 上一層
currLevel-1;若是,沿 出邊 (outgoing edges) 把未造訪鄰居標為當層。 - 需要快速取得「某 source 頂點的出邊」→ 即 adjacency matrix 某一列 (row) 的非零元素 → CSR(
srcPtrs+dst)。 - Push 語意:每個 active 頂點把自己的 depth「推 (push)」給所有鄰居。建 BFS tree 時等於「父找子」(top-down)。
// Fig. 15.6 Vertex-centric PUSH (top-down), uses CSR (srcPtrs, dst)
__global__ void bfs_push_kernel(CSRGraph csr, unsigned int *level,
unsigned int *newVertexVisited,
unsigned int currLevel) {
unsigned int vertex = blockIdx.x * blockDim.x + threadIdx.x; // 03
if (vertex < csr.numVertices) { // 04
if (level[vertex] == currLevel - 1) { // 05 在上一層才工作
for (unsigned int e = csr.srcPtrs[vertex]; // 06
e < csr.srcPtrs[vertex + 1]; ++e) { // 07 走出邊
unsigned int neighbor = csr.dst[e]; // 08
if (level[neighbor] == UINT_MAX) { // 09 未造訪?
level[neighbor] = currLevel; // 10 標當層
*newVertexVisited = 1; // 11 idempotent 旗標
}
}
}
}
}
Push: thread = vertex, 沿 row 掃出邊 (CSR)
level[] 0 1 1 X X X X X X (X = UINT_MAX 未造訪)
v1 active ──srcPtrs[1..2]──► dst: 3,4 ⇒ level[3]=level[4]=2
v2 active ──srcPtrs[2..3]──► dst: 5,6,7 ⇒ level[5..7]=2
只有「上一層」的 thread (v1,v2) 真正迭代鄰居,其餘 thread 立即退出
多個 thread 可能同時把同一鄰居標成同一個 currLevel、把 flag 寫成 1。因為大家寫的值完全相同,最終結果不變 → 這種 idempotent 寫入不是 read-modify-write,不需要 atomic。(在 15-Graph-Traversal/03-Frontiers-Privatization-and-Optimizations 用 frontier 時就必須改用 atomicCAS,因為還要把頂點推入 frontier,重複加入會造成浪費。)
Vertex-centric Pull(bottom-up / CSC)
- 每個 thread 負責一個 vertex,先檢查自己是否 尚未造訪;若是,沿 入邊 (incoming edges) 找是否有任何鄰居在 上一層;找到一個即把自己標為當層並 break。
- 需要快速取得「某 destination 頂點的入邊」→ adjacency matrix 某一行 (column) 的非零元素 → CSC(
dstPtrs+src)。 - Pull 語意:每個未造訪頂點「拉 (pull)」回前驅的 active 狀態。建 BFS tree 時等於「子找父」(bottom-up)。
// Fig. 15.8 Vertex-centric PULL (bottom-up), uses CSC (dstPtrs, src)
__global__ void bfs_pull_kernel(CSCGraph csc, unsigned int *level,
unsigned int *newVertexVisited,
unsigned int currLevel) {
unsigned int vertex = blockIdx.x * blockDim.x + threadIdx.x; // 03
if (vertex < csc.numVertices) { // 04
if (level[vertex] == UINT_MAX) { // 05 只處理未造訪頂點
for (unsigned int e = csc.dstPtrs[vertex]; // 06
e < csc.dstPtrs[vertex + 1]; ++e) { // 07 走入邊
unsigned int neighbor = csc.src[e]; // 08
if (level[neighbor] == currLevel - 1) { // 09 前驅在上一層?
level[vertex] = currLevel; // 10 標自己
*newVertexVisited = 1; // 11
break; // 12 找到一個就夠
}
}
}
}
}
只要一個前驅在上一層,就足以確定本頂點屬於當層,不必檢查其餘鄰居。只有「沒有任何前驅在上一層」的頂點(圖例中的 vertex 8)才會跑完整個鄰居清單。對 high-degree / 高變異的圖,提早 break 可大幅降低 load imbalance 與 control divergence。
Push vs Pull 與 Direction-Optimized
兩者都是 1 thread/vertex,但有兩個影響效能的關鍵差異:
| 差異點 | Push (top-down) | Pull (bottom-up) |
|---|---|---|
| 迴圈是否可早退 | 一定跑完整個鄰居清單 | 可 break 早退 |
| 哪些 thread 會迭代鄰居 | 只有上一層頂點 | 所有未造訪頂點 |
| 早期 level(前層小、未造訪多) | 較快(迭代少數清單) | 慢(很多未造訪頂點都要掃) |
| 後期 level(前層大、未造訪少) | 慢 | 較快(命中前驅機率高、易早退) |
早期 level 用 push、後期 level 切換成 pull。切換時機依圖型而定:
- 低 degree 圖(道路網):層數多、要很久才出現「大層 + 多數已造訪」→ 切換較晚。
- 高 degree / small-world 圖(社群網路):層數少、層成長快 → 切換較早。
Direction-optimized 需要 push 的 CSR 與 pull 的 CSC 兩種表示。
例外:若圖是 undirected(adjacency matrix 對稱,如社群網路、maze routing),CSR 與 CSC 等價,只需存一份即可供兩種 kernel 使用。
push=scatter(主動把值推給鄰居,寫入分散) ┐
pull =gather (主動把鄰居值拉進自己,寫入集中)┘ 與 Ch.17/18 的 scatter-vs-gather 同源
Edge-centric(COO)
- 每個 thread 負責一條 edge:用 COO
src[e]找源點、若源點在上一層,再用dst[e]找目的點、若目的點未造訪則標為當層。 - 需要由「一條邊」直接取得 source 與 destination(即某非零的 row index 與 column index)→ COO(
src+dst)。
// Fig. 15.10 Edge-centric, uses COO (src, dst)
__global__ void bfs_edge_kernel(COOGraph coo, unsigned int *level,
unsigned int *newVertexVisited,
unsigned int currLevel) {
unsigned int edge = blockIdx.x * blockDim.x + threadIdx.x; // 03
if (edge < coo.numEdges) { // 04
unsigned int vertex = coo.src[edge]; // 05 邊的源點
if (level[vertex] == currLevel - 1) { // 06 源點在上一層?
unsigned int neighbor = coo.dst[edge]; // 07 邊的目的點
if (level[neighbor] == UINT_MAX) { // 08 未造訪?
level[neighbor] = currLevel; // 09 標當層
*newVertexVisited = 1; // 10
}
}
}
}
Edge-centric: thread = edge (每 thread 恰好 1 條邊 → 無 load imbalance)
edges: (0,1)(0,2)(1,3)(1,4)(2,5)(2,6)(2,7)(3,4)(3,8)...
▲ ▲ ▲ ▲ ▲ ▲
src 在上一層的邊才寫入 dst;其餘 thread 檢查後即退出
優點 vs 缺點
| 優點 | 缺點 |
|---|---|
| 暴露更多平行度:邊通常遠多於頂點,小圖也能塞滿 device | 每條邊都檢查:頂點 v 即使無關,其 n 條邊各 launch 一個 thread 重複檢查 v n 次 |
| load imbalance / divergence 最低:每 thread 恰 1 邊(重排 thread→work 映射以減少 divergence,見 Ch.6) | COO 儲存成本高:需顯式存 src+dst 兩陣列,多於 CSR/CSC |
- vertex-centric 的對照:一個頂點 v 無關時,只啟一個 thread 檢查 v 一次即可 跳過全部 n 條邊。
小圖(頂點不足以塞滿 device)或 degree 變異極大(celebrity 頂點)的圖。
線性代數 / GraphBLAS 形式 (Linear-Algebraic Formulation)
- 上述 kernel 與 Ch.14 的 SpMV (sparse matrix-vector multiplication) 高度相似。稍作改寫,可把「一層 BFS 迭代」整個表達成 SpMV + 少數向量運算,其中 SpMV 是主導運算。
- 更廣義:許多圖計算都能以 adjacency matrix 的稀疏矩陣運算表達 → 稱為圖問題的 linear-algebraic formulation,由 API 規格 GraphBLAS 標準化。
| 優點 | 缺點 | |
|---|---|---|
| Linear-algebraic / GraphBLAS | 可直接複用成熟、高度優化的稀疏線性代數平行函式庫 | 錯失 針對特定圖演算法性質的最佳化(如 pull 的早退、frontier 跳過無關頂點) |
關鍵公式與複雜度 (Key Formulas)
- 滿連通有向圖邊數:
E_full = N(N-1)(無自環);稀疏圖中平均出度≪ N-1。 - 每層 launch 的 thread 數:
- Push / Pull(vertex-centric):
threads = V(全部頂點) - Edge-centric:
threads = E(全部邊;通常E ≫ V→ 平行度更高)
- Push / Pull(vertex-centric):
- 每層實際迭代鄰居的 thread 數:
- Push:
= |前一層頂點集合|(少) - Pull:
= |未造訪頂點集合|(早期多、後期少)
- Push:
- 每層工作量:vertex-centric 須掃所有頂點與相關邊
O(V + E);edge-centric 每層固定掃所有邊O(E)。 - COO 索引儲存:
2·E(src+dst) vs CSR/CSC 的(V+1) + E。
考試/面試重點 (Exam / Test Patterns)
| 情境 / 關鍵字 | 答案 / 技巧 |
|---|---|
| Push 用哪種格式?為何? | CSR:需「某 source 的出邊」= row 的非零元素 |
| Pull 用哪種格式?為何? | CSC:需「某 destination 的入邊」= column 的非零元素 |
| Edge-centric 用哪種格式? | COO:由一條邊直接取得 (src, dst) |
| Push 為何不需 atomic? | idempotent:所有 thread 寫相同值,非 read-modify-write |
Pull 為何能 break? |
只要一個前驅在上一層就足以判定,省去其餘檢查;降低 divergence |
| Push 用幾個 thread / 幾個迭代鄰居? | launch V 個;只有 上一層頂點 迭代鄰居 |
| Pull 用幾個 thread / 幾個迭代鄰居? | launch V 個;所有未造訪頂點 迭代鄰居 |
| Edge-centric launch 幾個 thread? | E 個(每邊一個),暴露平行度最高 |
| 早期 level 選哪個?後期? | 早期 push、後期 pull → direction-optimized |
| 何時切換點較早? | 高 degree / small-world 圖(層少、成長快)→ 切換早 |
| Direction-optimized 要存什麼? | 同時存 CSR + CSC;除非圖 undirected(對稱)只需一份 |
| 小圖 / degree 變異大選哪個? | edge-centric(平行度高、每 thread 1 邊無 imbalance) |
| Edge-centric 兩缺點? | 掃每條邊(無法跳過無關頂點的整串邊);COO 儲存多 |
| BFS 可用什麼線性代數表達? | SpMV + 少數向量運算;標準 API 為 GraphBLAS |
| GraphBLAS 缺點? | 錯失針對特定圖演算法的最佳化 |
| 為何每層各 launch 一個 kernel? | 需跨層 barrier:等上一層標完才標下一層,否則 level 標錯 |
| Push/Pull 對應 scatter/gather? | push=scatter(推給鄰居),pull=gather(拉自鄰居) |
Related Notes
- 15-Graph-Traversal/01-Graph-Representations-and-BFS
- 15-Graph-Traversal/03-Frontiers-Privatization-and-Optimizations
- 14-Sparse-Matrix-Computation/01-Sparse-Matrices-and-Basic-Formats-COO-CSR
- 18-Electrostatic-Potential-Map/01-DCS-Scatter-vs-Gather
- 17-Iterative-MRI-Reconstruction/02-FHD-Kernel-Parallelism-Structure
- 04-Compute-Architecture-And-Scheduling/02-Warps-SIMD-and-Control-Divergence