DeepSeek 开源 DeepGEMM-Ascend：支持华为昇腾 950 的矩阵乘法内核库，自测最高达硬件上限的 99.8%

9 月 30 日，DeepSeek 在 GitHub 的 deepseek-ai 组织下开源了新仓库，把用于大模型计算的通用矩阵乘法（GEMM）内核库 DeepGEMM 移植到华为昇腾 NPU（神经网络处理器），初版支持昇腾 950 设备。仓库 README 的更新记录写着：「2026.09.30：DeepGEMM Ascend 初始版本发布，支持 Ascend 950 设备；内核的设计目标是达到接近峰值的硬件性能。」

README 的性能表显示，稠密 GEMM 在多种数据类型上最高达到硬件上限的 99.8%。性能数字、兼容性与支持范围均为仓库自述口径。

初版以一条根提交发布：89 个文件、16,468 行新增

公开仓库的起点是 2026 年 9 月 30 日的一次「OSS release」提交（）：它没有父提交，是该公开仓库的根提交，包含 89 个文件变更、16,468 行新增、0 行删除；目录里放着 C++ 内核源码、Python 包、测试与性能分析脚本。

上游的仓库也在同一天的更新记录里写明 DeepGEMM Ascend 已经可用，并链向新仓库。

同一个包名、同一套 API：BF16、FP8、FP4 与 MoE 算子

README 首先强调兼容性：DeepGEMM-Ascend「完全兼容 DeepGEMM 的 API」，支持 BF16、FP8、FP4 GEMM、MQA logits 与 MegaMoE 算子（MoE 即混合专家）。它与 DeepGEMM 使用相同的包名，用户只需安装这个包，就能沿用其他平台上 DeepGEMM 的 API 和开发流程；内核接口文档直接引用上游 DeepGEMM 的接口说明。

上游 DeepGEMM 的硬件要求是 NVIDIA SM90 或 SM100 架构 GPU；这版移植把目标换成了华为昇腾 NPU。

移植覆盖的算子对应 DeepSeek 模型系列里的具体部件：MQA logits 是 DeepSeek Lightning Indexer 的评分算子；MegaMoE 把 EP dispatch、两个 grouped GEMM、SwiGLU 与 combine 融进一个内核；性能表里还包括面向 mHC（Manifold-Constrained Hyper-Connections）模块的 prenorm GEMM。

自测数据与测量条件：稠密 GEMM、MQA、MegaMoE 与 mHC

测量条件写得很具体：设备是 Ascend 950DT，工具链是 CANN 9.20，测量工具是 bench_msprof，缓存状态为 cold L2；形状与用例沿用 DeepGEMM 的测试套件，覆盖 DeepSeek 模型系列的推理与训练负载，README 称结果可在 tests/ 中复现。

稠密 GEMM 的四行结果都以 4,096×7,168×16,384 为形状：FP4×FP4 为 1,701 TFLOPS（每秒一万亿次浮点运算）、上限 1,730，利用率 98.3%；FP8×FP4 与 FP8×FP8 都是 861 TFLOPS、上限 865，利用率 99.5%；BF16×BF16 为 431 TFLOPS、上限 432，利用率 99.8%，对应延迟 2,229.9 微秒。

其余表格里：推理用 FP8 GEMM 覆盖 M=128 与 M=4,096 两组形状；MQA logits 内核受 FIX 流水线限制，利用率 99%；MegaMoE 在 EP8、每个 token 路由到 6 个专家、外加一个共享专家的配置下，最高测得 846.3 TFLOPS，数值跨 8 个 rank 取平均；mHC 的 prenorm GEMM 被描述为几乎使 HBM（高带宽内存）写带宽饱和，表格中最高的一行是 3,463 GB/s。README 的描述是：代码轻量，但在多种矩阵形状上能达到硬件极限性能。

MAD 抽象、稀疏加载与 int16 打包的缩放因子

为了在昇腾上让 GEMM 内核保持简洁高效，仓库对昇腾的矩阵乘加原语（MAD）做了一层轻量抽象，隐藏分形矩阵布局、对齐约束、地址计算与冗长底层参数带来的大部分复杂度；它还大量使用昇腾特有的优化手段，比如稀疏数据加载与基于协程的流水线，以接近硬件的性能极限。README 称这些实现也可以作为昇腾平台极致性能优化的参考。

README 用提示框特别提示了缩放因子（scaling factor）格式的差异：昇腾上的格式与 NVIDIA 不同，沿 K 维度每两个 UE8M0 缩放因子被打包进一个 int16，并按 MN-major 顺序存储；库里因此提供了 transform_sf_into_required_layout 等工具函数做布局转换。

工具清单还包括：set_num_sms / get_num_sms 控制内核可用的 AI core 数量；set_npu_arch / get_npu_arch 覆盖或查询 JIT（即时编译）所用的 dav-* 架构，传 0 则查询设备；ACLNN 参考 GEMM 用于在测试中交叉校验内核结果。JIT 运行时来自 DeepSeek 的，mHC 内核基于。

需求清单里的边界：昇腾 950 系列、CANN 9.20、MIT 许可

README 的需求清单列得很具体：华为昇腾 NPU（开发与验证都在昇腾 950 系列上完成）、提供 bin/bisheng 与 bin/ld.lld 的 CANN 9.20 工具链、Torch NPU 包 torch_npu、Python 3.10 或更高版本、支持 C++20 <format> 的编译器与标准库，以及被 HC prenorm 内核使用的 tilelang。安装方式是先带上 submodule 克隆仓库，再执行 pip install . --no-build-isolation。

仓库采用 MIT 许可（）；README 列出的项目负责人是 Kexing Zhou、Zhean Xu 与 Chenggang Zhao。

README 还致谢华为：「感谢华为在整个 DeepGEMM-Ascend 开发过程中提供的技术支持与工程经验。」