DeepSeek发布一批基于tilelang开发的TileKernels! 为V4预热?

链接：https://github.com/deepseek-ai/TileKernels/tree/main

这是昨天发布的，这批基于TileLang DSL编写的算子（Kernels）挺有意思，以下是具体的 Kernel 包含内容、特点，以及对 DeepSeek 新一代模型架构特点的个人推测：

仓库包含的核心 Kernel 类型

根据仓库代码，这批 Kernel 主要分为以下几个大类：

1. MHC，包含 pre_split_mixes、sinkhorn_normalize、pre_apply_mix 等算子。这是一个全新的网络拓扑结构算子集合，包含基于 Sinkhorn 算法的归一化和多头残差混合。

2. Engram Gate ，包含前向与反向融合算子 engram_gate_fwd/bwd

，将 RMSNorm、点积、带符号的平方根（signed_sqrt）以及门控聚合进行了极致的融合，支持权重梯度的 in-place 累加。

以上俩都在年初发布出来过论文，感兴趣可自行搜索阅读，仓库中开源了它们的tilelang实现

1. Quantization系列算子，包含极低比特量化算子，不仅有 FP8/FP4，还包含了一个罕见的自定义 12-bit 格式（E5M6）的 cast_e5m6，支持 per-token、per-block 和 per-channel 量化。甚至包含了前向 / 反向传播与激活函数的融合算子（如 swiglu_backward_and_per_token_cast），最近量化做的少了，12bit这个是新东西吗？反正对我来说挺有意思的，把quantization kernel和它之前的act+mul融合这个我到时以前见过，这不算是新东西

   
   

2. MoE Routing (混合专家路由)，包含 

topk_gate、reduce_fused、expand_to_fused 以及针对张量并行（TP）掩码优化的 Token 路由算子。

c. Transpose, 高度优化的批量矩阵转置算子，这个也不算新

这些 Kernel 的显著特点

• 
  

• 但是这些kernel全部由 TileLang 编写，这个算新。

• 
  

• 仓库主页明确提到这些kernel达到了硬件极限，不过没提哪个硬件:)

• 
  

• 比如以上的一些算子融合 (Kernel Fusion)方案为了突破访存瓶颈，DeepSeek 将大量连续操作写进同一个 Kernel 里。例如将激活函数、反向传播与张量 Cast（如 FP8/E5M6 转换）融为一体。

• 
  

• 面向训练，仓库中的算子大量包含带有 bwd 后缀的反向传播实现，并且在 modeling 目录下提供了完整的 torch.autograd.Function 封装。这说明这些 Kernel前后向一把抓

• 
  

猜测DeepSeek 新模型架构

💡特点一：抛弃传统残差连接， MHC板上钉钉（貌似说了句废话。。

代码中 mhc_mult = 4 和大量 mhc 算子表明，新模型放弃了 Transformer 标准的单一残差流。估计会维护了多个并行的流形残差（例如 4 个 Head 的残差流）。在每一层 Attention 或 FFN 之前，不同通道的信息会通过 sinkhorn_normalize（确保路由分配的平衡）和 mix 算子进行复杂的重组与交叉，经过计算后再聚合。这能极大地提升网络深层的信息流动能力和表征容量。

💡 特点二：引入 Engram模块（貌似也说了句废话。。

算子 engram_gate 实现了这样一个公式：

猜测新模型在传统的 Attention 和 FFN 之外，加入了一个专门的连续型键值记忆检索模块 (Key-Value Memory) 或新型的高级门控连接。它通过计算当前状态 x 与某种记忆键 k 的相似度，来决定注入多少记忆内容 v 到主干网络中。这可能是 DeepSeek 针对长文本记忆或世界知识强化设计的新型外挂。

💡 特点三：E5M6 (12-bit) 训练 / 推理数据格式

目前业界主要在卷 FP8 (8-bit) 和 FP4 (4-bit) ，E5M6（1 符号位 + 5 指数位 + 6 尾数位 = 12-bit）格式我个人还没见过，如果有朋友见过，还请评论留言。

猜测新模型由于 FP8 在某些高精度敏感算子（如梯度累加或某些高频特征）会掉点，而 BF16 又太耗显存了，DeepSeek 可能探索出了 12-bit 的甜点区格式。由于硬件通常以 8/16 字节对齐，源码中巧妙地将 3 个 E5M6 的元素（3×12=36 bits -> 结合对齐实际使用 96 bits / 8 elements = 12 bit/el）打包到一起（cast_e5m6.py），这暗示新模型在训练阶段可能就会部分激活这种非标准低比特格式，实现显存与精度的完美平衡。

💡 特点四：全面且极致的稀疏化 (MoE)

top2_sum_gate 等算子表明，新架构继续坚定不移地走 Sparse MoE 路线（每次路由选取 Top-2 专家等），并已经将 Token 到专家的分发 / 聚合（expand_to_fused）下沉到了最低层。这说明新一代模型仍是一个拥有海量参数但极度稀疏的玩意儿。

总结

DeepSeek 上传的这批 TileLang 算子，完全是为了高度定制化的自研模型架构铺路的。新模型在保留甚至强化 MoE 的基础上，肯定会在MHC、Engram以及E5M6 12-bit三个维度上做出创新，当然，肯定还有别的创新，我相信这不可能是所有的算子，肯定还有没有上传的，敬请期待吧。同时，依靠 TileLang本身就已经支持了 AscendC和AscendNPU IR backends，以及现在政策导向，极大概率还真是在昇腾上Day 0适配，让我们后面几天保持高度关注吧。