摘要： Cohere 於 2026 年 9 月 8 日發布 North Mini Code 的 megakernel serving engine，把 decode forward pass 合併為單一 persistent kernel。BF16、單張 H100 下，官方稱端到端比 vLLM 快 1.25 到 1.41 倍且量不到精度損失。實作是單一 CUDA 檔，最適合低 batch size、受記憶體頻寬限制的 decode 工作負載。
典型的 LLM serving stack 會把一次 forward pass 拆成好幾顆依序執行的 kernel：QKV 先 launch、然後等；attention 再 launch、再等；接著 MoE launch 完還是得等。單獨看每一次 launch 都稱不上問題，真正的開銷藏在兩次之間的空窗。到了 decode 階段、batch size 又不大的時候，GPU 有大半時間只是空等，真正花在運算上的比例並不高。
2026 年 9 月 8 日，Cohere 發布了為 North Mini Code 設計的 serving engine，主體是一顆 decode megakernel：BF16、單張 H100，端到端比 vLLM 快 1.25 到 1.41 倍。對產品開發者來說，重點不是「又多一招加速技巧」，而是它把「kernel 邊界」這項隱形成本攤在檯面上，同時示範了怎麼用單一個 CUDA 檔，做出真實 serving 需要的那些功能。
為什麼 decode 慢：不是算力不夠，是頻寬沒用滿
Autoregressive decoding 在低 batch size 下是 memory-bound 的。North Mini Code 是 30B 模型，每個 token 只有 3.3B 參數處於活躍狀態，BF16 之下每一個 decode step 都要從 HBM 搬 6.6 GB 的權重，8K context 另外還要加上約 0.5 GB 的 KV cache。H100 的 HBM 頻寬為 3.35 TB/s，換算出來的理論上限（Speed-of-Light）約 470 tok/s。但 vLLM 實際上只跑到 185 tok/s，等於頻寬僅用掉 39%。
Cohere 的 megakernel 在 batch size 1 時跑出 292 tok/s，也就是 62% 的 SoL，比 vLLM 快 1.58 倍。這段落差在不同的 batch size、以及最高 256K context 的範圍內都成立，而且量不到任何精度損失。
Megakernel 的運作方式：一個 persistent kernel 取代上百顆小 kernel
GPU 上有 100 到 150 個 SM（streaming multiprocessor），每個 SM 執行同一份 kernel 程式，處理的資料各不相同。Megakernel 的作法是：每個 SM 只 launch 一個 threadblock，在整個 decode step 期間維持 resident。工作分派不再由 driver 負責，改成每個 block 自己從 global memory 讀取一份 task list，每一個 task 對應某個 tile 的某一項 operation。
資料相依性也不再靠 kernel 邊界來表達，而是改用 global memory 裡的 counter：task 做完就 increment 對應的 counter，需要某份輸入時就 spin 等待那個 counter。如此一來，排程的顆粒度從「整個 operation」縮到「單一 tile」，同步的範圍也從「整顆 GPU」縮到「特定的 producer」。
三個實際的加速來源
Cohere 列出三個比「減少 launch overhead」更重要的因素，並依影響大小排序：
1. 消除 wave quantization
假設某顆 kernel 有 200 個 tile，而 GPU 有 132 個 SM。第一波先平行跑 132 個 tile，剩下的 68 個 tile 落到第二波，這個時候有 64 個 SM 是閒置的。kernel 切得越小，這種「進位損失」就越明顯。GEMM 的 tile 形狀受限於矩陣維度與 kernel 設計，tile 總數很少剛好是 SM 數的整數倍。Megakernel 沒有需要對齊的邊界：哪個 SM 有空，ready 的 tile 就直接在它上面開跑。
North Mini Code 在這裡特別吃香，原因是它採用 parallel transformer layer：attention 與 MoE feed-forward 都由同一份 normalized input 算出，最後再用 fused residual add + RMSNorm 合併。這代表 attention 與 MoE 之間不必互等輸出，megakernel 也就更能積極地把 ready 的工作「回填」給閒置的 SM。
2. 移除 false dependency
就算 workload 完全相同，各 SM 也不會同時把工作做完。kernel 邊界等同於一次全 GPU barrier，進度由最慢的那個 SM 決定。舉例來說，把 attention 拆成 4 個 KV group，其中某個 group 先算完，負責它的 SM 還是得等另外三個，儘管它下一步要用的資料早就在 memory 裡。Megakernel 改用 fine-grained barrier 消除這種假相依：某個 KV group 的 O-proj，可以在該 group 的 attention output 一落地就開始；MoE 的 down projection 也能在某個 expert 的 up projection 完成後隨即啟動，不必等所有 expert 到齊。
3. 權重 prefetch
權重是 immutable 的，不依賴這一步的 activation。因此 task 可以在 activation dependency 尚未滿足時，就先把權重 tile 從 HBM 搬進 shared memory。Cohere 用得最積極的地方是 router 與 QKV projection：在前一層 O-proj 的尾聲就開始 prefetch 權重，比 RMSNorm 的執行還早，把原本閒置的頻寬拿去用。
從 Hazy Research 的基礎出發，但做了務實的取捨
Cohere 的設計大量參考 Hazy Research 那篇開創性的「Look Ma, No Bubbles!」。該文把 Llama-3.2-1B 的 forward pass 融進單一 kernel，在 batch size 1 下達到 H100 頻寬的 78%，而 vLLM 與 SGLang 大約只有一半。Cohere 採用了其中三個關鍵想法：在 GPU 上實作 task interpreter pattern、counter-based synchronization，以及跨 task boundary 的 overlap。
不過 Cohere 也做了幾個不一樣的選擇。他們的 GEMM 大量使用 tensor core instruction（wgmma），即使在 batch size 1 也比 CUDA core 快，同時降低 register pressure。更關鍵的是，他們放棄用 shared memory paging 來做 weight prefetch——原本的設想是讓 memory load 在 previous task 釋放 buffer 之前就能開始，但實際上的 bookkeeping 太過複雜、容易產生 bug，帶來的 overhead 反而高過好處。
替代方案是兩個成本更低的重疊來源：一是在同類型的連續 GEMM task 之間，讓 pipeline 的 stage phase 跨 tile 延續下去，而不是每到一個 tile 邊界就 drain 再 refill；二是把重疊做進 GEMM pipeline 內部，讓 producer warp 在等待 cross-SM activation 之前，就先把 weight tile load 發出去。
對產品開發者的意義
這篇的價值不只是「Cohere 跑得比 vLLM 快」。它示範的是一種不同的思考方式：decode 階段的瓶頸不在 flops，而在 memory bandwidth 與 latency。如果你正在評估 serving 方案，與其緊盯著 benchmark 數字，不如先問：這個系統在 batch size 1 到 8 的區間裡，實際吃掉了多少 HBM 頻寬？kernel 邊界造成的等待又佔了多少比例？
Cohere 也強調，megakernel 並沒有想像中那麼難寫。他們的實作就是單一 CUDA 檔，沒有 compiler、沒有新的程式設計典範，只是把既有的 tiled GEMM 與 paged attention 重新組織成一套單一的 calling convention。對於想自己動手優化 decode 的團隊而言，這是一個務實的參考點。
如果你正在處理類似的效能問題，這篇的完整程式碼在 GitHub 上，可以直接看他們怎麼處理 task scheduler、barrier 和 weight prefetch。另外，如果你對「用更少的資源跑更大的模型」這個方向有興趣，之前寫過的細模型的大優勢：企業如何用小語言模型省錢又高效也討論了類似的取捨邏輯。
Megakernel 不是萬靈丹。它最適合的場景是 decode 階段、低 batch size、memory-bound。如果你的 workload 屬於 prefill-heavy 或 batch size 很大，傳統的 kernel-per-operation 可能還是比較合適。但當服務瓶頸落在 decode latency，而 batch size 多半很小的時候，把 kernel 邊界拆掉是值得認真評估的方向。