TileKernels 与 DeepGEMM 的分界线：DeepSeek 开源内核栈的分工

两个仓库不是重复造轮子，而是沿着一条清晰的技术边界切开的：

凡是能落到 Tensor Core 上的，进 DeepGEMM；凡是卡在访存带宽上的、或者还在改的新结构，进 TileKernels。

本文逐条对照它们的导出 API，标出各自独占的能力、真正重叠的三处，以及一层完整 MoE 前反向里每一步的归属。

仓库	语言	规模（排除 third-party）
DeepGEMM	CUDA / CuTe	26,472 行
TileKernels	TileLang / Python	13,996 行
DeepEP	CUDA + NCCL Gin	23,390 行

一、分工

DeepGEMM：冻结的算力核心

计算密集，算术强度高，目标是打满 Tensor Core（H800 上 1550 TFLOPS）
指令层面直接写 wgmma（SM90）和 tcgen05（SM100）的 PTX
算子集合基本固定：GEMM 的各种形状变体，多年值得手工调
只做 SM90 / SM100，靠 heuristics 为每代硬件单独选 tile

TileKernels：流动的结构前沿

访存密集，算术强度 $\le 1$ FLOP/byte，瓶颈在合并访存和 bank conflict
跑在 CUDA Core 的 SIMT ALU 上：位运算取指数、warp shuffle 归约、原子直方图
装的是还在演进的新结构：Engram、Manifold HyperConnection
TileLang 让一个融合 kernel 从「几周同步工程」变成「100 行 Python」

这条边界不是推断出来的，DeepGEMM 的 README 自己写死了：

Please note that operations like input transposition or FP8 casting must be handled separately by the user, please implement or fuse them into prior kernels independently. … our primary focus is on optimizing the GEMM kernels themselves.

—— DeepGEMM/README.md:72

「交给用户自己实现」—— TileKernels 就是那个「用户」。

三条判据

判据	DeepGEMM	TileKernels	证据
一	计算密集	访存密集	DeepGEMM 用 7 级软件流水 + warp 特化把 TMA 加载和矩阵计算叠起来；TileKernels 的 `group_count`、`batched_transpose`、`inplace_unique_group_indices` 全是带宽题
二	Tensor Core	SIMT ALU	DeepGEMM 只收能改写成 Tensor Core GEMM 的活——连注意力打分核 `sm90_fp8_mqa_logits.cuh` 都被硬掰成了 GEMM 形状
三	稳定	实验	GEMM 十年不变，值得手写 PTX；mHC 的 `pre_big_fuse` 把 RMSNorm + TF32 GEMM + split mixes + 10 轮 Sinkhorn 融进一发，用 CUDA 写要数周

二、TileKernels 有、DeepGEMM 没有

算子族	TileKernels API	DeepGEMM 有对应吗
Gating 门控打分	`topk_gate` `top2_sum_gate` `topk_sum_and_topk_group_idx` `normalize_weight` `aux_fi`	完全没有。DeepGEMM 自己的测试里直接用 `torch.topk` 顶着（`tests/test_mega_moe.py:125`）
MoE Routing token 重排	`get_fused_mapping` `expand_to_fused` / `_with_sf` `reduce_fused` `group_count` `mask_indices_by_tp` `inplace_unique_group_indices`	没有独立算子。只在 SM100 的 `fp8_fp4_mega_moe` 里被吞进大 kernel 内部，取不出来单用
Quantization 量化	`per_token_cast` / `per_block_cast` `per_channel_cast`（+transpose / fused） `per_token_cast_to_e5m6` `cast_back` / `cast_back_e5m6` `swiglu_forward_and_per_token_cast` `swiglu_backward_and_per_token_cast`	部分重叠。DeepGEMM 没有任何 BF16→FP8/FP4 的 cast kernel，只有 scale factor 的重排工具 `transform_sf_into_required_layout`。唯一例外是它 vendor 了一个 TileLang 写的 SwiGLU+cast
Transpose	`transpose` `batched_transpose`	没有。DeepGEMM 靠 `nt/nn/tn/tt` 四种 TMA 描述符隐式处理转置
Engram 远程记忆	`engram_hash` `engram_gate_fwd` / `bwd` `fused_weight` `grad_w_reduce`	零支持。Engram 的另一半在 DeepEP：取数走 0-SM RDMA（`impls/engram_fetch.cuh`），算 gate 走 TileKernels
mHC Manifold HyperConnection	`expand_to_mhc` `sinkhorn_normalize` `pre_split_mixes` / `pre_apply_mix` `post` / `head_compute_mix` `pre_big_fuse` `multilayer_recompute`	只重叠 1 个。`tf32_hc_prenorm_gemm` 对应 TileKernels 的 `_mhc_pre_norm_fn_fwd_mul`。Sinkhorn、head mixing、post 组合 DeepGEMM 一概没有
Modeling 可训练层	`modeling.engram.engram_gate` `modeling.mhc.functional`	没有。DeepGEMM 是纯 kernel runtime，不提供 `torch.autograd.Function` 抽象

三、DeepGEMM 有、TileKernels 没有

一句话：TileKernels 里没有任何通用 GEMM。唯一沾边的是 mHC 里那个不支持 split-k 的专用矩乘。

能力	API	用在哪
稠密 GEMM	`fp8_fp4_gemm_{nt,nn,tn,tt}` `bf16_gemm_{nt,nn,tn,tt}` `cublaslt_gemm_*`	所有 dense 层
M-grouped 连续	`m_grouped_fp8_fp4_gemm_nt_contiguous` `m_grouped_bf16_gemm_nn_contiguous`	MoE 前向 + 激活反向，每个专家 token 数不等且不补齐
M-grouped masked	`m_grouped_fp8_fp4_gemm_nt_masked`	解码期定形张量，配合 CUDA Graph；输入可直接吃 DeepEP 低延迟 kernel 的输出
K-grouped 连续	`k_grouped_fp8_gemm_tn_contiguous`	MoE 权重梯度 $dW = X^T dY$ ，所有专家一次 launch 算完
Lightning Indexer	`fp8_fp4_mqa_logits` `fp8_fp4_paged_mqa_logits` `fp8_gemm_nt_skip_head_mid`	DeepSeek-V3.2 的稀疏注意力打分；`skip_head_mid` 跳过中间 head 段写出，把 MLA 的解耦 RoPE 和压缩 latent 分开
Einsum	`einsum` / `fp8_einsum`	`bmk,bnk->mn`、`bhr,hdr->bhd` 等注意力投影收缩
Mega MoE	`fp8_fp4_mega_moe` `bf16_mega_moe`	SM100 单 kernel 融合 EP dispatch + L1 + SwiGLU + L2 + combine，通信走 symmetric memory
JIT 运行时	`csrc/jit/*`	C++20 运行时编译、TMA 描述符直建、PDL 开关。DeepEP 复用了同一套骨架

四、真正重叠的三处

A · SwiGLU 前向 + 量化

DeepGEMM 的 third-party/tilelang_ops/ 里 vendor 了一个 TileLang 写的 swiglu_apply_weight_to_fp8（212 行）。它不是 CUDA，就是 TileKernels 那一套写法搬过来的——目的是让 test_mega_moe.py 能跑非融合 baseline 而不用装外部包。

差别：

TileKernels 版支持 E4M3 / E5M2 / E2M1 / E5M6 全格式、per-token 和 per-channel 两种、还有配套反向
vendor 的这个只有前向、只有 FP8 E4M3
Mega MoE 内部那份根本不落地——在 L1 和 L2 之间的寄存器/共享内存里就算完了

B · mHC 的 pre-norm GEMM

两边都算 $D = A \cdot B^T$ 同时归约 $\sum A^2$ 给 RMSNorm 用。区别是 split-k，而且 TileKernels 的作者在注释里直接认了：

# TileLang implementation doesn't support split-k, so we set n_splits to 1
# You may want to adopt the DeepGEMM implementation with split-k for better performance
n_splits = 1

出自 TileKernels/tile_kernels/modeling/mhc/ops/norm_fn.py:87，同样的注释也出现在 ops/pre_big_fuse.py:50。

DeepGEMM 的 tf32_hc_prenorm_gemm 有跨 SM 的 grid 级 split-k（多开一维 [num_splits, M, N]），token 数少时占用率明显更好。这是唯一一处两边功能真重、且明确写着该用哪个的地方。

C · Scale factor 的排布

DeepGEMM 要求 scale factor 在 SM90 上列主序 + TMA 16 字节对齐，在 SM100 上打包成每个 torch.int 塞 4 个 UE8M0。

DeepGEMM 自己有 transform_sf_into_required_layout 做独立一趟转换
TileKernels 则在量化/重排的时候就直接按这个物理布局写出去，省掉一次 kernel launch——它的 use_tma_aligned_col_major_sf 和 use_packed_ue8m0 两个 flag 就是为此存在的

这说明 TileKernels 不是「碰巧配合」，是照着 DeepGEMM 的 TMA 消费端反推着设计的生产端。

五、代码里互相引用的地方

除了上面 mHC 那两条注释，还有一条透露了路线图：

# All kernels may be deprecated in the future (or rewrite in TileLang)

出自 DeepGEMM/deep_gemm/legacy/__init__.py:1——这个 legacy 目录装的是 A100 的 Triton kernel。

也就是说 DeepSeek 的方向是：非 GEMM 的遗留 Triton 算子，系统性地重写成 TileLang。TileKernels 不是 DeepGEMM 的补丁，是它主动让出去的那一半。

六、一层 MoE 的完整前反向

步骤编号是真实执行顺序。

#	阶段	做什么	算子	来自
1	FWD	门控打分，选 top-2 组再选 top-k 专家	`top2_sum_gate`	TileKernels
2	FWD	算 token→专家映射表、各专家起止、补齐	`get_fused_mapping`	TileKernels
3	FWD	输入激活量化到 FP8 并产出 scale factor	`per_token_cast`	TileKernels
4	FWD	按专家段重排 token 和 sf，顺带排成 TMA 列主序	`expand_to_fused_with_sf`	TileKernels
5	FWD	L1 GEMM $X_{exp} \cdot W_1^T$	`m_grouped_fp8_gemm_nt_contiguous`	DeepGEMM
6	FWD	SwiGLU 前向 + FP8 量化 + sf 打包，一发融合	`swiglu_forward_and_per_token_cast`	TileKernels
7	FWD	L2 GEMM $Inter_{exp} \cdot W_2^T$	`m_grouped_fp8_gemm_nt_contiguous`	DeepGEMM
8	FWD	按路由权重加权合回原 token 布局	`reduce_fused`	TileKernels
9	BWD	把 $dY$ 沿 top-k 展开到专家位置	`expand_to_fused`	TileKernels
10	BWD	$dY_{exp}$ 量化到 FP8	`per_token_cast`	TileKernels
11	BWD	L2 激活梯度 $dY_{exp} \cdot W_2$	`m_grouped_fp8_gemm_nn_contiguous`	DeepGEMM
12	BWD	L2 权重梯度 $dW_2 = dY_{exp}^T \cdot Inter_{exp}$	`k_grouped_fp8_gemm_tn_contiguous`	DeepGEMM
13	BWD	SwiGLU 反向 + FP8 量化	`swiglu_backward_and_per_token_cast`	TileKernels
14	BWD	L1 激活梯度 $dGateUp_{exp} \cdot W_1$	`m_grouped_fp8_gemm_nn_contiguous`	DeepGEMM
15	BWD	L1 权重梯度 $dW_1 = dGateUp_{exp}^T \cdot X_{exp}$	`k_grouped_fp8_gemm_tn_contiguous`	DeepGEMM
16	BWD	专家激活梯度累加回输入梯度 $dX$	`reduce_fused`	TileKernels

数一下：16 步里 DeepGEMM 只占 6 步，全是 GEMM；其余 10 步全在 TileKernels，全是重排、量化、门控。

这就是「分工」最直观的样子——DeepGEMM 负责算，TileKernels 负责把数据摆成能算的形状再摆回来。

如果是 SM100 上的推理，fp8_fp4_mega_moe 可以把第 4–8 步整体替换成一个 kernel，但第 1–3 步仍然得由 TileKernels 做。