FlashQLA：Qwen 把線性注意力推進 DGX Spark 的 2-3× 加速 kernel

2026-04-29，Qwen 團隊開源 FlashQLA——基於 TileLang 的 GDN (Gated Delta Network) linear attention kernel library，forward pass 比 FLA Triton kernel 快 2–3×、backward 快 2×。它的硬體支援列表很有意思：SM90+（Hopper、H100、H200），以及 SM121 的 GB10 系統，也就是 NVIDIA DGX Spark。

這篇拆三件事：①FlashQLA 解的是什麼問題、為什麼 2-3× 重要；②TileLang 是什麼、為什麼 Qwen 不用 Triton；③FlashQLA 在 Qwen 整個技術棧的位置——它不只是優化，是 Qwen 把線性注意力推進 personal device 的關鍵一步。

FlashQLA 解的是什麼

要看懂這個 kernel 的重要性，得先知道 GDN 是 Qwen 的什麼。

從 2025 末的 Qwen3-Next 開始，GDN 成為整個 Qwen 家族的主力 attention 層——Qwen3-Next-80B-A3B 一路到 Qwen3.5 / Qwen3.6 系列（397A17B / 122A10B / 35B / 27B）。Qwen 的設計是 hybrid：GDN 跟標準 full attention 交替使用，目標是同時拿到「full attention 的表達力」和「linear attention 在長序列的效率」。

GDN 是線性注意力（linear attention），把 softmax attention 的 O(n²) 複雜度降到 O(n)。差別在 GDN 用一個帶指數衰減的 gate 控制歷史 context 怎麼被帶進來——這個 gate 既是它跟標準 linear attention 的差別，也是 FlashQLA 用來變快的關鍵。

問題在哪：當模型擴大到 35B、122A10B、397A17B，context 拉到 256K 以上，GDN 在訓練與推理的端到端開銷已經變成不能忽略的瓶頸。原本社群在用的 FLA Triton kernel 在 Hopper 上沒有把硬體吃滿——Tensor Core、CUDA Core、SFU（special function unit）三類運算單元各有閒置。

FlashQLA 的目標就是把這個閒置吃掉。

三個技術設計

照官方 README + NYU Shanghai RITS 的解析，FlashQLA 的 2-3× 來自三個合在一起的設計：

1. Gate-driven 自動 intra-card CP

利用 GDN gate 的指數衰減特性，在 TP（Tensor Parallelism）、長序列、小 head count 三類場景自動啟用 intra-card context parallelism。意思是：當單一 TP shard 沒辦法把整張 GPU 餵飽時，FlashQLA 自動把 context 維度的工作再切一刀，讓多個 SM 同時處理同一 shard 的不同 context 片段。

這個設計直接對應到 edge inference 場景——personal device 上你只有一張 GPU，模型放上去後 batch 通常很小，不切 CP 的話 SM 大量閒置。

2. Hardware-friendly algebraic reformulation

把 GDN Chunked Prefill 的 forward + backward 數學重寫，降低三類 GPU 單元的負擔：

• Tensor Core：負責矩陣乘法
• CUDA Core：負責 scalar / vector 運算
• SFU：負責 exp、sqrt 這類超越函數

關鍵是「重寫但不損失數值精度」——這對 GDN 特別重要，因為它的 recurrent state 會放大 rounding error，kernel 偷工減料的話訓練會發散。

3. TileLang fused warp-specialized kernels（拆成兩個）

這是論文裡明說的反直覺取捨：FlashQLA 沒有把整個 GDN flow 融合成一個 kernel，而是拆成兩個——一個針對 CP 優化，一個針對 backward 效率優化。

代價：在大 batch 上會多出 kernel 之間的 memory I/O 開銷。 回報：在 edge device + 長 context 場景，整體 wall-clock 反而更快。

backward 那一支是最難的：他們建了一條 16-stage warp-specialized pipeline，在極緊的 on-chip memory 限制下塞進去，最終拿到 2×+ 的 kernel-level 加速。

TileLang 是什麼、為什麼不用 Triton

過去 LLM kernel 的事實標準是 OpenAI 的 Triton。FlashQLA 用 TileLang 不是隨意的選擇。

TileLang 是 2025 年 ICLR 論文（arxiv.org/pdf/2504.17577）發表的 GPU kernel DSL，核心概念是把 tile（一塊形狀化的資料區塊）當成 first-class 物件。它對開發者暴露的就是 tile 級的抽象——你直接寫「這塊資料從 global 搬到 shared memory，用 Tensor Core 算 GEMM，結果寫回」，編譯器自己處理 warp specialization、async copy、pipeline。

論文裡的數字相當激進：

• H100 上 vs Triton：Flash MLA 達 4.06–10.59× 加速，且程式複雜度比 hand-tuned FlashMLA 少 6.86×
• AMD MI300X 上 vs Triton：5.64–12.97× 加速，略勝 hand-tuned ROCm AITER 函式庫 (1.05×)
• 一個 kernel 的程式碼量「從幾百行降到 80 行以下」

對 Qwen 來說，TileLang 的優勢是兩個：

1. 能榨出 Hopper 的特定指令（cp.async.bulk、wgmma.mma_async），這些 Triton 不容易直接表達
2. 開發成本低——backward 的 16-stage pipeline 用 Triton 寫會痛苦得多

Qwen 賭 TileLang 而非 Triton，本質是賭「在 Hopper 上把硬體吃乾」這件事 Triton 不夠用。

FlashQLA 為什麼挑這個時機

時機點是這篇文章最該說的東西。

過去半年的 NVIDIA 個人 AI 硬體有兩件大事：H200 普及到二級雲、DGX Spark（GB10 SM121）開始出貨。DGX Spark 是 NVIDIA 把 Grace + Blackwell 塞進桌上電腦的產品，目標是「讓 AI 工程師在自己機器上跑大模型」。

但 DGX Spark 一個著名的弱點是 LPDDR5X 記憶體頻寬——decode 階段被頻寬卡死，dense 模型實測常常掉到個位數 tok/s。社群的應對是用 MoE 模型（每 token 只活化一小部分參數）來繞開。

Qwen3.6 / Qwen3.5 的架構正好是 MoE + GDN hybrid——MoE 解 decode 的記憶體頻寬問題，GDN 解 prefill 跟長 context 的計算問題。FlashQLA 是這套架構的 prefill 加速器。FlashQLA 的硬體支援列表寫死「SM90+ 包含 SM121 GB10」，意圖很清楚——他們要 Qwen 在 DGX Spark 上跑得起來、跑得好。

換句話說：FlashQLA 不是孤立的 kernel 優化，是 Qwen 端到端策略的關鍵一塊。模型架構（MoE+GDN）+ 推理 kernel（FlashQLA）+ 硬體選型（DGX Spark）= Qwen 在 personal device 上跑 agentic AI 的完整解法。

edge agentic AI 為什麼需要這個

The Rift 對這次發布的觀察值得引用：「2-3× forward speedup is material for real-time agent decision cycles」。這個 framing 對。

Agent 跟普通 chatbot 的差別在於 agent 一次任務通常要做 多次 LLM call + tool call——一個 task 跑下來可能 20–50 次 round trip。每次 round trip 的 latency 都在累積。雲端 agent 你看不見這個 cost（被分散到 OpenAI / Anthropic 的帳單），personal device 上每一毫秒都是你自己的 GPU 時間。

把 prefill（agent 每次處理新 prompt 的階段）加速 2-3×，等於 agent 整體任務時間少 30-50%。對「agent 跑半小時 vs 跑十分鐘」這種等級的任務，這是質的差別。

這也解釋了為什麼 Qwen 在推文裡明確寫「Purpose-built for agentic AI on your personal devices」——不是 generic 性能優化，是策略性的 edge agent 鋪路。

一個容易忽略的細節

FLA Triton kernel 是社群現在的事實標準（Songlin Yang 主導的 flash-linear-attention 專案）。Qwen 把比較對手放成 FLA 而不是某個自家 baseline，是個刻意的選擇——它在告訴社群「我們不是在跟自己比，我們在跟你們現在用的東西比」。

NYU Shanghai 的 RITS 報告裡還補了一個數字：FlashQLA 對 FlashInfer 0.6.9 在同樣 shape 也大約 2× 加速。FlashInfer 是 NVIDIA / Apache 主導的 LLM inference kernel 集合，這意味著 FlashQLA 不只贏了社群方案，也贏了大廠方案。MIT License 釋出意味著任何用 GDN 架構的模型——不只 Qwen，將來其他基於 Gated DeltaNet / KDA 的模型——都可以受惠。

一句話收尾

GPU kernel 的發布通常不會上熱門推文。FlashQLA 上得了，是因為它代表的不是又一個「快了 2 倍」的微優化，是 Qwen 把整套「MoE 架構 + 線性注意力 + 個人裝置」的押注，往執行層完成了關鍵一塊拼圖。下一次你在 DGX Spark 上跑 Qwen3.6 的時候，FlashQLA 就在底下默默讓 prefill 不卡。