GLM 推論導入 EPD 分離架構，吞吐衝到 3 倍
Zhipu AI（z.ai）Blog · 2026-09-17
GLM-5.3-Flash 的推論服務把 Prefill、Decode 與多模態 Encode 拆成三個獨立階段，改用 Encode-Prefill-Decode（EPD）分離式架構取代原本單體式的推論流程，再疊上張量並行、W8A8 量化、混合精度 KV cache 與 ReplaySSM，把端到端服務吞吐拉高到基線的三倍。這套系統由智譜（Zhipu AI）從零打造在超過十萬顆國產 AI 加速卡上，扛下 GLM-5.3-Flash 的全部生產流量，細節寫在 z.ai 於 2026-09-17 發布的部落格裡。
原本的問題
這是第一次有團隊在十萬顆規模的國產加速卡叢集上部署生產推論，晶片記憶體容量與頻寬相對有限，同時還要撐住新模型架構、1M token 的上下文視窗與多模態請求。生態系不成熟、kernel 支援不完整，很多本該有文件的地方只能用猜的。整個適配到上線的過程，主要靠 GLM-5.3 驅動的 Infra Agent 完成，而不是純靠人力堆工程師。
Agent 驅動最佳化最大的瓶頸不是寫程式能力，而是回饋太粗：端到端指標只能告訴 agent「TTFT 漲了 30%」或「吞吐掉了 20%」，卻不會告訴它問題出在哪一層。團隊因此把正確性測試、執行追蹤、runtime event、microbenchmark 都接進 agent 的迭代流程，讓每個假設都能在本地驗證，不必每次改動都跑一次完整的服務部署與端到端壓測。
採用的方法
第一個案例是KDA kernel 的 Context Parallelism（CP）路徑算錯。CP 需要合併不同 context 分片的 state，核心運算為：
M = tl.dot(M_chunk, M) # 合併分片間的 state transformation
S_next = tl.dot(M, S) + H # 更新下一分片的初始 state
原本 tl.dot 在輸入是 FP32 時仍預設用 TF32 計算，長上下文下誤差會持續累積。修法是把兩處都明確設成 input_precision="tf32x3"，用三次 TF32 Tensor Core 運算合成出接近 FP32 的精度，同時盡量保留 TF32 的效能。這個修正已經合併進上游 Flash Linear Attention PR #1180。
第二個案例是KV Transfer 被 Python GIL 卡住。團隊訂的驗收標準是 Prefill + KV Transfer 跟純 Prefill 的效能差不能超過 5%，但 agent 量到部分場景差了 20% 以上。往下追時間軸發現，Python 端的 KV Transfer 從沒跟 DeepEP 的 dispatch／combine 重疊過；追進 DeepEP v1.2.1 的程式碼才發現 intranode_dispatch／intranode_combine 都沒有釋放 GIL，CPU 等 GPU 回傳 token 數時，同一行程裡負責 Mooncake Transfer 的 Python 執行緒也一起被鎖住，傳輸沒辦法跟運算重疊。同版本的 internode_dispatch 就有明確釋放 GIL，程式裡的註解正是為了不擋住其他執行緒的 KV Transfer，這個對照佐證了 agent 的判斷。修法是在那幾段 C++ 執行區間釋放 GIL，修完後 Prefill + KV Transfer 跟純 Prefill 的效能差降到 1% 以下。
實際效果
Kernel 層的最佳化同樣靠這套回饋機制滾動：以 KDA Decode kernel 為例，從 v0 到量產版本經過三輪迭代。
| 版本變化 | 改動 | 對執行時間的影響 |
|---|---|---|
| v0 → v1 | 導入 ReplaySSM，用計算換記憶體 | 執行時間不減反增 |
| v1 → v2 | Agent 做除法最佳化 | 比 v1 快 9.6% |
| v2 → 量產版 | 合併 V 維度分塊，把重複四次的 FP32 normalization／gating 換成單次 warp 級 reduction | 比 v2 快 1.71 倍 |
整套 EPD 分離架構加上量化與 ReplaySSM，讓 GLM-5.3-Flash 從第一次成功跑起來到正式上線只花不到兩週，端到端吞吐相對最初基線衝到三倍，硬體使用率與每 token 成本也追平主流 NVIDIA GPU。上線後，GLM-5.3-Flash 以匿名代號 Ox-Alpha 在 OpenCode 與 OpenRouter 上測試，一週內就成為兩個平台上最常被呼叫的模型，六天內處理超過 62 兆個 token。
對正在跑開源大模型 serving stack 的團隊，這裡有兩個可以直接比對的檢查點：用 Triton 寫 CP／TP 這類切分 kernel 時，要確認 tl.dot 等運算在 FP32 輸入下有沒有偷用 TF32，尤其長上下文場景誤差會被放大；另一個是用 DeepEP、Mooncake 這類 EP 通訊庫時，得檢查 C++ 端有沒有在等待 GPU 結果時持有 GIL，同行程內其他 Python 執行緒的排程可能因此被卡住，讓原本設計好的非同步重疊完全發揮不出來。