AI slob越來越進步 今天有點累 明天再mur mur
Day 22 的 Crossing evaluator 很老實。
如果 graph 有 E 條 edges,它讓每一個 GPU thread 負責一條 edge,再把其他 E 條 edges 全部掃過一次:
thread e
-> crossing(e, 0)
-> crossing(e, 1)
-> crossing(e, 2)
-> ...
-> crossing(e, E - 1)
-> output counts[e]
整體工作量是:
E * E
這個版本很適合當 correctness baseline,因為每一份 candidate layout 都從頭算,不需要相信任何舊狀態。
但 SA 一次通常只做這件事:
把一個 node v
從 old_position
移到 new_position
其他 nodes 沒有動,絕大多數 edges 也沒有動。
結果我們卻把整張圖的 edge pairs 又算了一次。
所以今天的優化不會先從 shared memory、warp shuffle 或 instruction unrolling 開始。我們先問一個更值錢的問題:
哪些 crossing 根本不可能改變?
假設 node v 連著三條 edges:
● a
/
/ e0
/
b ●---● v---● c
e1
\ e2
\
● d
e0、e1、e2 都接到 node v,因此它們叫做 v 的 incident edges。
我們把這個集合寫成:
I(v) = all edges connected to node v
集合大小就是 node 的 degree:
d(v) = size(I(v))
當 v 被移動時,只有 I(v) 裡面的線段會改變形狀:
before move
● a
/
/
/
● v
after move
● a
___/
● v'
沒有接到 v 的 edge,兩個端點都沒動,所以線段也完全相同。
於是我們得到今天最重要的 invariant:
if edge e is not in I(v)
and edge f is not in I(v):
crossing_old(e, f) == crossing_new(e, f)
這不是近似,也不是「通常不會改變」。只要兩條線段的四個端點都沒變,它們的 crossing 狀態就一定不變。
移動 node v 後,任意一對 edges 只會落在下面三種情況。
第一種,兩條都是 static edges:
e not in I(v)
f not in I(v)
兩條線段都沒動
-> crossing 不變
-> 不用重算
第二種,一條是 incident edge,另一條是 static edge:
e in I(v)
f not in I(v)
其中一條線段改變
-> crossing 可能出現
-> crossing 也可能消失
-> 必須比較 old 與 new
第三種,兩條都是 v 的 incident edges:
e in I(v)
f in I(v)
兩條 edges 共用 node v
-> 不會被算成 proper crossing
第三種在 crossing count 中不會增加數量,但 geometry validator 仍然要檢查移動後是否出現共線重疊等非法情況。Crossing evaluator 少算工作,不代表合法性檢查可以一起刪掉。
把三組放在一起看:
partner edge f
incident static
+------------+------------+
incident e | shared v | may change |
+------------+------------+
static e | may change | unchanged |
+------------+------------+
右下角的 static-static 區域通常最大,也是我們真正要刪掉的工作。
E * E 變成 d(v) * EDay 22 的 full evaluator 不管移動哪個 node,都做:
work_full = E * E
增量更新只需要處理 incident edges 和其他 edges 的關係,主要工作量變成:
work_delta is proportional to d(v) * E
假設 graph 有 512 條 edges,而這次移動的 node degree 是 8:
full scale
512 * 512
= 262,144
affected-pair scale
8 * 512
= 4,096
這裡的 4,096 是用來理解數量級的 affected-pair scale,不是 kernel 實際執行的精確 instruction count。實作還需要比較 old/new crossing,並為 incident 與 static edges 各自產生正確的 per-edge count。
但方向已經很清楚:
d(v) << E
時,
d(v) * E << E * E
一般 graph 的 node degree 往往遠小於總 edge 數,因此演算法上少算,通常比把完整雙迴圈微調快一點更有價值。
先看一條沒有接到 moved node 的 static edge j。
它和其他 static edges 的 crossing 狀態全部不變,所以不必從零開始計數。新的 count 可以從舊值修正:
new_count[j] = old_count[j]
接著只檢查它與每一條 incident edge i 的 old/new 狀態:
delta(i, j) = new_cross(i, j) - old_cross(i, j)
其中 crossing Boolean 可以當成 0 或 1:
old = 0, new = 0 -> delta = 0
old = 0, new = 1 -> delta = +1
old = 1, new = 0 -> delta = -1
old = 1, new = 1 -> delta = 0
最後:
new_count[j] = old_count[j]
+ sum(
new_cross(i, j) - old_cross(i, j)
for i in I(v)
)
這裡同時要算 old 與 new,不能只看 candidate layout。
如果一條 crossing 因為 node 移開而消失,只看 new layout 只會得到 0,卻不知道應該從舊 count 減掉 1。
再看一條接到 moved node 的 incident edge i。
它本身的幾何位置已經改變,理論上它與任何 static edge 的 crossing 都可能不同。因此最清楚的做法是讓它從 0 開始,掃過所有 partner edges:
new_count[i] = 0
for every partner edge p:
new_count[i] += new_cross(i, p)
這個 branch 仍然需要 E 等級的工作,但只有 d(v) 條 incident edges 會進來。
完整邏輯是:
for each output edge j:
if j is incident to moved node:
recount j against all partner edges
else:
begin with old_count[j]
apply old/new delta from I(v)
我們仍然維持 Day 22 的 ownership:
一個 thread
-> 負責一條 output edge
-> 使用 local accumulator
-> 最後寫一次 output
因此即使變成增量更新,也不需要讓許多 threads 對同一個 per-edge counter 做 atomicAdd。
如果每個 proposal 都掃過全部 E 條 edges,詢問「這條 edge 有沒有接到 node v」,我們只是把成本藏到另一個迴圈裡。
Graph topology 在整個搜尋期間不會改變,所以 node-to-edge adjacency 應該只建立一次。
原專案使用 CSR 形式:
node_edge_offsets[V + 1]
incident_edges[2 * E]
例如:
edge 0 = (0, 1)
edge 1 = (0, 2)
edge 2 = (1, 3)
edge 3 = (2, 3)
每個 node 的 incident list 是:
node 0 -> [0, 1]
node 1 -> [0, 2]
node 2 -> [1, 3]
node 3 -> [2, 3]
壓平成兩個 arrays:
offsets = [0, 2, 4, 6, 8]
incident = [
0, 1,
0, 2,
1, 3,
2, 3
]
要查 node 2,只需要:
begin = offsets[2] = 4
end = offsets[3] = 6
incident[4:6] = [1, 3]
在 CUDA kernel 裡就是:
int begin = node_edge_offsets[moved_node];
int end = node_edge_offsets[moved_node + 1];
for (int k = begin; k < end; ++k) {
int edge_id = incident_edges[k];
// process affected edge
}
為什麼 incident_edges 大小是 2 * E?因為一般 edge 有兩個端點,會分別出現在兩個 nodes 的 incident lists 裡。
這個索引在 CPU 建立一次,再上傳 GPU 一次。之後每個 proposal 都重用,不需要在 hot loop 裡重新配置或重新傳輸。
build once
-> upload once
-> reuse for millions of proposals
這個優化介紹的不是更快的乘法,而是資料結構如何決定 GPU 最後要做多少工作。
假設 graph 有 N 個 nodes,我們同時評估 P 個 candidate moves。
最直接的資料表示可能是替每個 candidate 複製一份完整座標:
candidate_positions[P][N][2]
但一個 candidate 只移動一個 node。其餘 N - 1 組座標和 base layout 完全相同。
所以範例只保存三個整數:
move[p] = {
node_id,
new_x,
new_y
}
讀取座標時才疊上 candidate:
int ax = (a == moved_node) ? new_x : base_x[a];
int ay = (a == moved_node) ? new_y : base_y[a];
如果有 128 個 nodes、32 個 proposals,使用 32-bit integer:
複製完整 layouts
32 * 128 * 2 * 4 bytes
= 32,768 bytes
只保存 moves
32 * 3 * 4 bytes
= 384 bytes
每批 proposals 額外需要的座標資料從 32 KB 降到 384 bytes;base layout 仍然只在 GPU 保存一份。graph 越大,差距越明顯。
這就是很典型的 GPU memory optimization:不是想辦法讓 PCIe 搬得更快,而是不要搬那些從來沒有改變的資料。
教學範例同時評估 32 個 independent proposals。
Grid 可以想成一張二維表:
edge index
0 1 2 3 ... E-1
proposal 0 [ ] [ ] [ ] [ ] ... [ ]
proposal 1 [ ] [ ] [ ] [ ] ... [ ]
proposal 2 [ ] [ ] [ ] [ ] ... [ ]
...
proposal P-1 [ ] [ ] [ ] [ ] ... [ ]
CUDA mapping 是:
int p = blockIdx.y;
int e = blockIdx.x * blockDim.x + threadIdx.x;
每個 cell 負責:
proposal p 的 edge e 新 crossing count
核心 kernel 可以縮成:
if (edge e is incident to moved_node) {
count = 0;
for (int f = 0; f < num_edges; ++f) {
count += crossing_after_move(e, f);
}
} else {
count = base_count[e];
for (int f : incident_edges[moved_node]) {
count += crossing_after_move(e, f)
- crossing_before_move(e, f);
}
}
output[p * num_edges + e] = count;
當 E 很小時,單一 proposal 可能只有一兩個 blocks,根本填不滿 GPU。把多個 independent proposals 放進 blockIdx.y,可以提供更多 blocks,增加 throughput。
但這裡有一條不能偷換的語意:
32 個 independent proposals from one base layout
不等於
32 個 sequential SA steps
真正的 SA 如果接受 proposal 0,proposal 1 理論上應該從更新後的 layout 出發。教學 benchmark 中的 32 個 proposals 只是用來測試 evaluator throughput 與 correctness,沒有 commit moves,也沒有模擬一條完整 SA chain。
同一個 warp 裡,大部分 edges 通常是 static,少數 edges 是 incident:
lane 0 -> static branch
lane 1 -> static branch
lane 2 -> incident branch
lane 3 -> static branch
...
incident branch 會掃 E 條 partners;static branch 只掃 d(v) 條 incident edges。兩邊工作量不一樣,因此 warp 可能需要等待較慢的 branch 完成。
這是增量演算法的代價之一:工作變少了,但規則不再像 full evaluator 那麼整齊。
進一步的實作可以把工作拆成不同 kernels:
kernel A
更新 static candidate edges
kernel B
重新計算 incident edges
這樣每個 kernel 裡的 threads 做比較一致的工作,但會多一次 kernel launch,也需要額外的 candidate list。是否划算,要看 graph 大小、degree 分布與 proposal batch size。
所以「避免 divergence」也不能脫離整個 pipeline 單獨判斷。
原專案把前面的 dense incremental scaffold 稱為 C1。
C1 已經把 pair tests 從 E * E 降到 d(v) * E 等級,但每個 proposal 仍會為所有 E 條 output edges 啟動 worker:
incident edge
-> full recount
non-incident edge
-> scan I(v) and apply delta
如果 graph 很大,下一個問題是:連 E 條 static edges 都有必要逐條看嗎?
答案仍然是否定的。
一條 incident segment 從舊位置移到新位置,只可能與空間上靠近它舊路徑或新路徑的 edges 改變 crossing 關係。這就是 C3 spatial candidate evaluator 的出發點。
先把畫布切成格子,並記錄每一格有哪些 edges 經過:
+-------+-------+-------+-------+
| | e7 | | |
+-------+-------+-------+-------+
| e2 | e2,e7 | e7 | |
+-------+-------+-------+-------+
| | e5 | e5 | e5 |
+-------+-------+-------+-------+
當 node v 移動時,對每一條 incident edge 取得:
old segment bounding box
new segment bounding box
再查詢兩者覆蓋的 grid cells:
old cells
union
new cells
-> candidate partner edges
-> deduplicate
-> exact crossing test
只有 candidate edges 需要計算 old/new delta。其他 non-incident、non-candidate edges 可以直接複製舊 count:
new_count[j] = old_count[j]
Spatial grid 只負責縮小 candidate set,最後的答案仍然由 exact integer crossing predicate 決定。它不是用格子近似 crossing。
很多人第一次寫 spatial delta,會只查新線段經過的 cells:
new bounding box
-> 找現在可能 crossing 的 edges
這只能抓到新產生的 crossing,卻可能漏掉已經消失的 crossing。
例如:
before move
incident edge i X static edge j
old_cross(i, j) = 1
after move
incident edge i static edge j
new_cross(i, j) = 0
如果 j 只出現在 old segment 的區域,查詢 new region 時根本看不到它。counts[j] 就不會減掉 1,錯誤會留在 state 裡。
因此 candidate set 必須來自:
candidates = edges_in(old_region)
union
edges_in(new_region)
這是增量演算法很常見的陷阱:新狀態告訴你增加了什麼,舊狀態才告訴你失去了什麼。
Spatial grid 通常使用固定容量,因為我們不想在每個 proposal 裡動態配置記憶體。
但固定容量一定會遇到 overflow:
某個 cell 裡的 edges 太多
或
union 後的 candidate edges 太多
最危險的做法是只保留前 capacity 條 edges:
candidate_count > capacity
-> silently truncate
這會讓 kernel 很快,答案卻不再 exact。
原專案的正確策略是:
overflow detected
-> 記錄 typed fallback reason
-> 放棄這次不完整的 candidate list
-> 回到 dense C1 或 full-square path
-> 只發布 exact result
fallback 偶爾比較慢沒有關係。對 optimizer 來說,罕見的慢答案仍然比無聲的錯答案好;錯誤 crossing counts 一旦被 accepted move 寫回 current state,後面的 SA 會在錯誤狀態上繼續跑,而且可能很久才被發現。
C3 需要標記哪些 edges 是本次 candidates。最直接的方式是:
is_candidate[E] = all zeros
for edge in candidates:
is_candidate[edge] = 1
但這代表每個 proposal 都要先清空整個 E 長度的 array。
另一種方法是使用 generation stamp:
current_generation = 137
candidate_generation[edge] = 137
判斷時比較:
candidate_generation[edge] == current_generation
下一次 proposal 只要把 generation 加一,不必先把整個 array 歸零:
proposal 137 -> generation 137
proposal 138 -> generation 138
proposal 139 -> generation 139
這個技巧省掉的是固定的 memory write。當我們已經努力把 crossing tests 減少後,這類每回合都掃完整 array 的初始化成本就會開始變得顯眼。
當 generation integer 即將 wrap 時,仍需要安全地 reset;這也是實作不能省略的邊界條件。
教學範例使用 RTX 4060 Laptop GPU,固定 seed = 42,一次 kernel 評估 32 個 independent proposals。
結果如下:
| Graph | Full median | Dense delta median | Full / Delta |
|---|---|---|---|
| 48 nodes / 96 edges | 0.02473 ms | 0.02243 ms | 1.10x |
| 128 nodes / 512 edges | 0.25788 ms | 0.13190 ms | 1.96x |
這裡的 delta 是本篇前半段的 dense incremental C1 教學 kernel,沒有包含 C3 spatial grid。
測量範圍是:
CUDA events
5 次 warm-up
9 batches
每個 batch 連續執行 100 次
最後取 median
包含的是 GPU stream 上的 kernel 執行區間。不包含:
CuPy JIT compile
memory allocation
host-device transfer
Python CPU oracle
完整 V1~V4 validator
SA acceptance 與 commit
512-edge case 的 crossing work 在演算法上減少很多,實測卻只有大約 1.96x。這其實很合理:
第一,C1 仍然為每條 output edge 啟動 thread
第二,static branch 要同時計算 old 與 new crossing
第三,CSR incident list 帶來額外的間接讀取
第四,incident 與 static branch 工作量不同,可能產生 divergence
第五,output 仍然是 proposals * edges 個 counts
第六,小 kernel 仍受固定 launch 與排程成本影響
96-edge case 只快 1.10x,也不是優化失敗。工作太小時,省下來的幾何測試還不足以壓過固定成本。
這是本篇最值得保留的 GPU 經驗:
演算法複雜度下降
不代表
小資料一定立刻得到同等比例的 speedup
資料量、parallelism、memory access 與固定成本,最後都會反映在 wall time 裡。
Full recompute 每次從座標重新建立 counts。Delta evaluator 則依賴:
current coordinates
current per-edge counts
candidate move
如果某次更新漏掉一條 edge,錯誤會被寫進 current counts。下一次 proposal 又以錯誤 counts 當 base,偏差可能一路累積。
因此測試不能只看最後的 K,必須逐 edge 比較:
delta_counts(candidate)
==
full_recompute(candidate)
至少要覆蓋:
移動後新增一個 crossing
移動後移除一個 crossing
crossing pair 的兩條 edges 都正確更新
完全不受影響的 edge count 保持不變
degree 0 node
degree 1 node
high-degree node
重複接受多次 moves 後仍然沒有 drift
spatial old region 能抓到消失的 crossing
candidate overflow 正確 fallback
教學範例會讓 GPU full counts、GPU delta counts 與獨立 CPU oracle 逐 edge 相等,才接受結果。
範例中的 proposals 都從同一份 immutable base layout 出發,因此很容易驗證:不需要把任何 candidate 寫回 current state。
真正的 SA 則必須分成兩個階段:
evaluate
把 candidate coordinates 與 counts 寫到 scratch
計算 candidate objective
accept
才把 coordinates、counts、objective 一起 commit
拒絕 proposal 時:
current coordinates unchanged
current counts unchanged
current objective unchanged
接受 proposal 時,這三份狀態必須一起前進:
new coordinates
new per-edge counts
new K, n_K, Phi, C
不能讓其他 kernel 看見「新座標配舊 counts」,也不能先發布新的 objective、稍後才補 crossing state。在單一 block 內可以用 cooperative phases 與 synchronization;跨 kernels 則要靠同一 stream 的 ordering 保證提交順序。
這部分不會出現在單獨的 crossing microbenchmark 裡,卻是把演算法放進完整 GPU solver 時最容易踩到的工程問題之一。
本篇 benchmark 只量 crossing-count kernel。
得到新的 counts[e] 之後,solver 還要算:
K
n_K
Phi
C
最直接的方法是對完整 E 長度的 scratch counts 做 reduction。這已經比 E * E 的 full crossing evaluator 便宜很多,但仍然是一筆成本。
更深的版本可以維護 crossing-count histogram,只更新真正改變的 buckets,再從 non-empty buckets 找新的 K。代價是 state 與 commit 邏輯會更複雜。
因此這篇的 1.96x 只能回答:
dense full crossing-count kernel
和
dense incremental crossing-count kernel
哪一個比較快?
它不能回答完整 SA solver 最後快了幾倍。要得到那個答案,必須把 validation、objective reduction、proposal generation、acceptance、commit 與同步全部放進相同的 measurement scope。
Day 23 做的不是把 Day 22 kernel 換成比較花俏的 CUDA 語法,而是重新定義需要執行的工作:
(node, new_x, new_y),不複製完整 layout。Day 22 的問題是:
怎麼把所有 crossing counts 算對?
Day 23 的問題則是:
既然只移動一個 node,
怎麼保留不變的答案,只更新真的可能改變的部分?
下一篇會進入 Layout。Crossing evaluator 已經告訴我們一份 layout 有多好,接下來要問的是:第一份 node positions 從哪裡來?為什麼隨機亂放雖然簡單,卻可能讓後面的 SA 一開始就背著大量 crossings?