圖走訪與廣度優先搜尋 練習題 (Practice - Graph Representations and Breadth-First Search)

Fully connected 有向圖共 N(N−1) 條邊（無自環，每點 out-degree = N−1）。

• CSR(srcPtrs+dst)：易取某 source 的出邊（row 非零）→ push / top-down。
• CSC(dstPtrs+src)：易取某 destination 的入邊（column 非零）→ pull / bottom-up（也用於最短路徑回溯找 predecessor）。
• COO(src+dst)：由一條 edge 直接取 (src, dst) → edge-centric。
  核心觀念：格式選擇 = 演算法選擇。

label = 從 root 到該 vertex 的最少邊數（最短 hop count = level/depth）。
每層各 launch 一個 kernel 是因為需要一個跨層 barrier：必須等上一層所有 vertex 都標完，才能開始標下一層，否則可能把某 vertex 標到錯誤的 level。換 root 後 BFS 結果可能截然不同（即使新 root 只差 1 條邊）。
複雜度 O(V + E)（每 vertex、每 edge 各走訪一次）。

thread = 一個 vertex。檢查 level[vertex] == currLevel − 1（我的 vertex 在上一層）；若成立，沿 CSR 出邊（srcPtrs/dst）把所有未訪問（level == UINT_MAX）鄰居標為 currLevel，並設 newVertexVisited = 1。
不需 atomic 因為這是 idempotent 寫入：多個 thread 即使同時標同一鄰居，寫的都是相同的值（currLevel、1），非 read-modify-write，最終結果不變。

使用 CSC(dstPtrs+src，入邊)。每個 thread 先檢查自己 level[vertex] == UINT_MAX（尚未訪問）；若是，沿入邊找是否有任一鄰居（前驅）在 currLevel − 1，找到就把自己標為 currLevel 並 break。
能 break 因為：只要一個前驅在上一層就足以判定本 vertex 屬當層，不必再檢查其餘鄰居。只有「沒有任何前驅在上一層」的 vertex 才會跑完整串清單。對 high-degree / 高變異的圖，提早 break 可大幅降低 load imbalance 與 control divergence。

使用 COO（src+dst）。
優點：(1) 暴露更多平行度 — 邊通常遠多於頂點（E ≫ V），小圖也能塞滿 device；(2) load imbalance / divergence 最低 — 每 thread 恰好處理 1 條邊（屬於重排 thread→work 映射以減 divergence，見 Ch.6）。
缺點：(1) 掃每一條邊 — 無關頂點 v 的 n 條邊各 launch 一個 thread，重複檢查 v 共 n 次（vertex-centric 只需 1 thread 檢查 v 一次即跳過全部 n 條邊）；(2) COO 儲存成本高 — 需顯式存 src+dst 共 2E，多於 CSR/CSC 的 (V+1)+E。

每層只 launch numPrevFrontier 個 thread（= 上一層 frontier 的大小），而非全圖 V 個 → 消除冗餘 thread。
必須用 atomicCAS 是因為 frontier 版除了標 level，還要把鄰居入隊到 currFrontier。若多個 thread 都看到該鄰居未訪問，就會把它重複入隊，導致下一層重複處理 → 浪費。atomicCAS(&level[neighbor], UINT_MAX, currLevel) 一次原子完成「檢查未訪問 + 標記」，回傳值 == UINT_MAX 者才是第一個訪問者，只有它能入隊。
結束條件：numCurrFrontier == 0（不再需要 newVertexVisited flag）。

競爭最嚴重的是全域 atomicAdd(numCurrFrontier, 1) — 所有成功訪問的 thread 都對同一個 global counter 累加。（atomicCAS 競爭只是中等，因為僅少數 thread 撞同一鄰居。）
Privatization：每個 block 在 shared memory 維護自己的 local frontier 與 local counter，thread 只與同 block 的 thread 競爭（shared atomic 低延遲）；commit 時每個 block 只對 global counter 做 1 次 atomicAdd 配置整段空間。
額外好處：local → global 的 commit 用 idx_s = threadIdx.x; idx_s += blockDim.x 步進，使連續 thread 寫連續 global 位址 → coalesced stores。（local frontier 溢位時需把 counter 釘回 LOCAL_FRONTIER_CAPACITY 並改走 global atomicAdd fallback。）

Launch overhead：小 frontier 時，終止舊 grid + 啟動新 grid 的開銷可能超過平行收益。解法：準備一個只用單一 thread block 的 kernel，以 __syncthreads() 在 level 之間同步、連續跑多個小 level；當 frontier 大到溢出 block-level frontier 時，把內容複製到 global frontier 並返回 host，host 再改用一般 multi-grid kernel。
Load imbalance（celebrity vertex degree 高出數個數量級）兩解法：(1) 改用 edge-centric；(2) 把 frontier 內 vertex 依 degree 分桶（Merrill & Garland）：small → 1 thread / vertex、medium → 1 warp / vertex、large → 1 block / vertex。

各 vertex 出邊：v0→{1,3}、v1→{2}、v2→{}、v3→{2,4}、v4→{1}，共 E = 6。

• srcPtrs = [0, 2, 3, 3, 5, 6]（長度 N+1 = 6；末元素 6 = 總邊數）。
  • v0 起點 0（2 邊）、v1 起點 2（1 邊）、v2 起點 3（0 邊，故 srcPtrs[2]==srcPtrs[3]）、v3 起點 3（2 邊）、v4 起點 5（1 邊）。
• dst = [1, 3, 2, 2, 4, 1]。
  驗證：v3 出邊 = dst[srcPtrs[3]..srcPtrs[4]−1] = dst[3..4] = {2, 4} ✓。

進入本輪前已標：0(L0)、1(L1)、3(L1)；未訪問 = {2, 4}。

• Push：launch V = 5 個 thread；只有 vertex 在 currLevel−1 = 1 的 thread 迭代鄰居 → {1, 3} 共 2 個。
• Pull：launch V = 5 個 thread；迭代鄰居的是所有未訪問頂點 {2, 4} 共 2 個；兩者都找到上一層前驅（2 的入邊有 1,3；4 的入邊有 3）→ 2 個 thread 標記自己。
• Edge-centric：launch E = 6 個 thread（每邊一個）；src 在 level 1 的邊 = 1→2, 3→2, 3→4 → 3 個 thread 可能 label vertex。

參數：addr = 資料記憶體位址、compare = 用來比較的值、val = 比較成功時要寫入的值。語意（原子）：if (*addr == compare) *addr = val;。回傳值 = *addr 的舊值。
用法：atomicCAS(&level[neighbor], UINT_MAX, currLevel) — 把 level[neighbor] 與 UINT_MAX（未訪問）比較，相等則設成 currLevel。
成功判定：若回傳值 == UINT_MAX，代表此 thread 是第一個把該鄰居從未訪問改成已訪問的 thread → 只有它負責把 neighbor 入隊，避免重複入隊。

差異 1（迴圈早退）：push 一定跑完整個鄰居清單；pull 找到一個前驅在上一層即 break，對 high-degree / 高變異圖能大幅降低 load imbalance 與 control divergence。
差異 2（哪些 thread 迭代鄰居）：push 只有「上一層頂點」迭代；pull 是「所有未訪問頂點」都迭代。
early level：上一層頂點少、未訪問頂點多 → push 只掃少數清單（快），pull 要掃大量未訪問頂點（慢）。
late level：上一層頂點多、未訪問頂點少，且命中已訪問前驅而早退的機率高 → pull 快，push 慢。
Direction-optimized = 早期 push + 後期 pull；high-degree / small-world 圖切換較早，low-degree 圖切換較晚。需同時存 CSR + CSC；若圖 undirected（adjacency matrix 對稱）則兩者等價，只需一份。
對應：push = scatter（主動把 depth 推給鄰居，寫入分散）、pull = gather（主動把鄰居狀態拉進自己，寫入集中），與 Ch.17/18 的 scatter-vs-gather 同源。

• 平行度：edge-centric = E（最高，E ≫ V，小圖也能塞滿 device）；vertex-centric = V（小圖可能 thread 不足）。
• load imbalance / divergence：edge-centric 最低（每 thread 恰 1 邊，等同 Ch.6 重排 thread→work 映射）；vertex-centric 依各 vertex 的 degree，高變異圖嚴重 imbalance。
• 儲存成本：edge-centric COO = 2E（高）；vertex-centric CSR/CSC = (V+1)+E（低）。
• 跳過無關工作：vertex-centric 能跳過 — 一個無關 vertex v 只需 1 thread 檢查一次就略過全部 n 條邊；edge-centric 不能 — v 的 n 條邊各 launch 一個 thread，重複檢查 v 共 n 次。
  適用：edge-centric → 小圖、或 degree 變異極大（celebrity）的圖；vertex-centric → degree 較均勻的圖，且可進一步搭配 frontier 剪掉無關 thread、再用 privatization 降低 atomic 競爭。

核心訊息：格式選擇 = 演算法選擇（CSR↔push、CSC↔pull、COO↔edge）；BFS = O(V+E) 逐層 wavefront；最佳化路徑 = vertex/edge-centric → direction-optimized → frontier（剪冗餘）→ privatization（降競爭）→ launch reduction + degree 分桶（降開銷、平衡負載）。