剛剛，DeepSeek 悶聲開源，觸摸 GPU 性能上限

DeepSeek 又悄悄開源了一個項目，一套直接榨干 H100/H200/B200 的 GPU Kernel 庫，名字叫 TileKernels

它和之前放出來的 FlashMLA、DeepGEMM、DeepEP 是一個級別的「內功心法」，只不過這次換了TileLang來寫

簡介

TileKernels 是 DeepSeek 開源的一組專門為 LLM 訓練和推理優化的 GPU Kernel 集合

用 DeepSeek 官方自己的話說：

? Most kernels in this project approach the limit of hardware performance regarding the compute intensity and memory bandwidth. Some of them have already been used in internal training and inference scenarios.

翻譯一下就是：這些 kernel 基本已經摸到硬件性能上限了，無論是算力還是顯存帶寬，而且其中相當一部分已經在 DeepSeek 內部的訓練和推理里用過了

換句話說，這是從 DeepSeek 自家流水線里拆出來直接開源的真家伙，不是實驗室玩具

核心功能一覽：

• Gating — MoE 路由里的 Top-k 專家選擇和打分

• MoE Routing — Token 到 Expert 的映射、融合擴展/歸約、權重歸一化

• Quantization — Per-token / per-block / per-channel 的 FP8 / FP4 / E5M6 量化，還把 SwiGLU 和量化融合在一起做

• Transpose — 批量轉置

• Engram — 融合了 RMSNorm、前向/反向傳播、權重梯度歸約的 Engram 門控 kernel

• Manifold HyperConnection (mHC) — 包含 Sinkhorn 歸一化、混合拆分/應用的超連接 kernel

• Modeling — 用 torch.autograd.Function 把底層 kernel 包成可訓練層（Engram Gate、mHC Pipeline）

看完這個列表我的第一反應是：這不就是 DeepSeek V3 / R1 那套 MoE + FP8 訓練體系的核心零件嗎？

• MoE 路由、Gating、Token-to-Expert 映射 → V3 的 MoE 架構

• FP8 / FP4 / E5M6 量化 + 融合 SwiGLU → V3 的低精度訓練和 V3.2 的量化推理

• Engram、Manifold HyperConnection → 這倆名字就很"研究院"，大概率和后續模型架構相關

這里插播一下，TileKernels 不是用傳統 CUDA 或 Triton 寫的，而是用的 TileLang（tile-ai/tilelang）

TileLang 是一門基于 TVM 的領域特定語言（DSL），專門用來寫高性能 GPU Kernel

它的定位和 Triton 有點像，但更"Pythonic"，而且官方號稱：

• 80 行 Python 代碼就能寫出 H100 上和 FlashMLA 打平的 MLA Decoding

• 支持 H100 (Auto TMA/WGMMA)、A100、MI300X、甚至華為昇騰和 Apple Metal

• 自動優化，寫得省心，跑得還快

所以 DeepSeek 選 TileLang 的意圖就很明顯了：用更簡潔的方式寫出性能逼近手工 CUDA 的 kernel。對于想學底層優化、又被 CUDA 勸退過的人，這套開源組合（TileLang + TileKernels）簡直是神級教材。

安裝

環境要求就一句話：你得有張 H100 級別的卡

• Python 3.10+

• PyTorch 2.10+

• TileLang 0.1.9+

• NVIDIA SM90 或 SM100 架構 GPU （也就是 H100 / H200 / B200 這個檔位）

• CUDA Toolkit 13.1+

裝法有兩種：

# 本地開發版（可改代碼）
pip install -e ".[dev]"


 # 直接裝發行版
pip install tile-kernels

這里就勸退一大波人了——SM90 起步，4090/A100 都不在支持列表里

測試與壓測

官方給了 pytest 跑法，可以只驗正確性，也可以跑 benchmark：

# 單文件：只驗正確性，4 個 worker 并行
pytest tests/transpose/test_transpose.py -n 4

 # 單文件：正確性 + 性能 benchmark
pytest tests/transpose/test_transpose.py --run-benchmark

 # 全量壓測（開 FULL_TEST 環境變量，跑 2 輪）
TK_FULL_TEST=1 pytest -n 4 --count 2

對做底層優化、學 kernel 的同學來說，這套測試腳手架本身就值得抄一份——正確性測試 + benchmark 一條龍，還帶壓力測試模式，工程完成度很高

目錄結構

tile_kernels/
├── moe/        # MoE 路由相關 kernel
├── quant/      # FP8/FP4/E5M6 量化
├── transpose/  # 批量轉置
├── engram/     # Engram 門控 kernel
├── mhc/        # Manifold HyperConnection kernel
├── modeling/   # 高層 autograd 封裝（Engram、mHC）
├── torch/      # PyTorch 參考實現（對標用）
└── testing/    # 測試和 benchmark 工具

最值得說的是 torch/ 這個目錄——每個 kernel 都配了一份純 PyTorch 的參考實現

這意味著你不僅能看到高性能版本，還能對著慢速但易懂的 PyTorch 版對照學習

這對想啃底層優化的同學來說，簡直就是"雙語對照教材"

我怎么看這個項目

先說優點，相當直接：

1. 真家伙，不是 PPT 。DeepSeek 自己內部訓練用過的 kernel，性能逼近硬件上限

2. 雙語對照 。每個 kernel 都配 PyTorch 參考實現，學習成本大幅降低

3. 工程質量高 。測試、benchmark、壓測腳手架全都配齊，MIT 協議隨便用

4. 覆蓋面精準 。MoE + FP8 量化這兩塊，恰好是當前大模型訓練最吃性能的地方

再說局限，也得真誠：

1. 硬件門檻高 。SM90 起步，家用卡基本別想，連 A100 都不支持。

2. 文檔幾乎為零 。官方自己也承認："they do not represent best practices and we are actively working on improving the code quality and documentation." 目前想吃透，只能啃代碼。

3. 依賴 TileLang 。這是一個相對新的 DSL，社區規模還沒 Triton 那么大，生態有待培養

4. 面向研究者 / 框架開發者 。如果你只是調調 API、煉煉小模型，這套東西你用不著；它是給寫訓練框架、寫推理引擎的人看的

DeepSeek 這次開源的 TileKernels，氣質和之前開源周放出來的 FlashMLA、DeepGEMM、DeepEP 一脈相承——不玩大新聞，就是把自己實打實用過的、能摸到硬件天花板的底層代碼直接公開

適合誰看：

• 想深入理解 MoE 訓練底層 kernel 怎么寫的同學

• 在做推理引擎、訓練框架、想榨干 H100 性能的工程師

• 對 FP8 / FP4 低精度訓練感興趣的研究者

• 想系統學習 TileLang 這門新 DSL 的朋友（配合 TileLang 官方的 Puzzle 學習路徑更香）

最后一句感慨：現在這個時代，開源一個模型權重已經算不上什么大新聞了，真正能體現一家公司技術深度的，是愿不愿意把底層這些"內功心法"也放出來

DeepSeek 在這件事上，誠意一直很足

制作不易，如果這篇文章覺得對你有用，可否點個關注。給我個三連擊：點贊、轉發和在看。若可以再給我加個，謝謝你看我的文章，我們下篇再見！