deepseek-ai/DeepGEMM-Ascend
DeepGEMM-Ascend: clean and efficient matrix multiplication kernel library for Huawei Ascend NPUs
DeepGEMM‑Ascend 是什么?
DeepGEMM‑Ascend 是开源 DeepGEMM 矩阵乘法库的硬件特定移植,目标是华为的 Ascend NPU 系列(例如 Ascend 950)。它实现与 DeepGEMM 相同的 Python 级 API,因此为其他平台(CUDA、ROCm 等)编写的代码只需安装此包即可在 Ascend 上运行。
此库以轻量 C++/Python 扩展包装 Ascend MAD(矩阵加法乘法)原语,处理分形布局、对齐和地址计算等低级细节。它还加入 Ascend 专属优化,例如稀疏数据加载和基于协程的流水线,目标是在各种矩阵形状下达到硬件的理论峰值性能。
谁可能使用它?
| 受众 | 重要性 |
|---|---|
| AI 研究人员/工程师构建在 Ascend NPU 上运行的大型语言模型(LLM)——他们需要快速的 GEMM、MQA 和 MoE 内核。 | |
框架开发者扩展 PyTorch‑NPU(torch_npu)以加入自定义高性能运算。 |
|
| 性能工程师寻找极致优化 Ascend 内核的参考实现。 | |
| 硬件感知的 ML 从业者想要在 Ascend 上基准测试 BF16、FP8、FP4 和混合精度矩阵运算。 |
核心功能(依 README 所述)
| 功能 | 详细说明 |
|---|---|
| API 兼容性 | 使用与上游 DeepGEMM 项目相同的包名和函数签名。 |
| 精度支持 | BF16、FP8、FP4 GEMM;也提供 MQA logits 和 MegaMoE 的专门内核。 |
| 轻量抽象 | 隐藏 Ascend 特定的布局和对齐复杂性,同时保持内核代码简洁。 |
| Ascend 特定优化 | 稀疏数据加载、基于协程的流水线、自定义缩放因子打包(UE8M0 配对 → int16)。 |
| 工具函数 | 设置 AI 核心数、选择 NPU 架构、控制确定性模式、转换缩放因子布局,以及查询理论对齐限制。 |
| JIT 运行时 | 依赖 DeepJIT 运行时;内核在首次使用时编译并缓存(可通过 DG_JIT_* 环境变量配置)。 |
| 环境变量调整 | 调试、缓存位置、启动超时,以及其他通过 DG_JIT_* 和标准 Ascend 变量暴露的设置。 |
| 性能 | 基准测试显示密集 GEMM 在 BF16/FP8/FP4 下可达硬件限制的 98‑99%,并提供分组 GEMM、MQA、MegaMoE 和自定义「HC Prenorm」内核的详细数据。 |
如何开始(快速入门摘要)
- 硬件与软件先决条件
- 华为 Ascend NPU(已在 950 系列验证)
- CANN 9.20 工具包(
bisheng、ld.lld) torch_npuPython 包- Python 3.10+、兼容 C++20 的编译器,以及
tilelang、tree‑sitter、tree‑sitter‑cpp
- 使用子模块克隆
git clone --recursive https://github.com/deepseek-ai/DeepGEMM-Ascend.git cd DeepGEMM-Ascend - 构建 C++ 扩展(仓库提供
develop.sh脚本,会链接 include 并运行编译器)。 - 安装
pip install . --no-build-isolation - 在 Python 代码中使用相同的 DeepGEMM API,例如
import deep_gemm as dg dg.set_num_sms(8) # 选择要使用的 AI 核心数 out = dg.aclnn_fp8_fp4_gemm_nt(A, B, scaling_factors)
何时表现最佳?
README 列出数个基准测试类别,其中库接近 Ascend 硬件上限:
- 密集 GEMM – BF16、FP8、FP4 可达理论 TFLOPS 的 99.8%。
- M‑Grouped GEMM – 对 Mixture‑of‑Experts(MoE)专家有用,在实际专家‑token 模式下可达 > 800 TFLOPS。
- MQA Logits – 多查询注意力的快速评分,延迟在 100‑300 µs 范围。
- MegaMoE – 端到端融合内核(dispatch + 两个 GEMM + SwiGLU),同时保持高计算和内存带宽。
- HC Prenorm GEMM – 针对「Manifold‑Constrained Hyper‑Connections」模块的专门内核,饱和 HBM 写入带宽。
这些数字显示此库已准备好用于 Ascend 硬件上的生产级 LLM 推理/训练工作负载。
限制/需注意事项
- 硬件特定 – 仅适用于华为 Ascend NPU;没有 CPU 或 GPU 后备方案。
- 工具链繁重 – 需要 CANN 9.20、
bisheng编译器,以及支持<format>的 C++20 编译器。 - 环境变量调整 – 性能可能取决于正确设置
DG_JIT_*变量和 Ascend 启动阻塞/调试标志。 - 测试版状态 – 仓库报告为初始版本(2026‑09‑30);社区支持和文档可能仍有限。
在 AI 生态系统中的定位
DeepGEMM‑Ascend 位于 AI 堆栈的低级线性代数层。现代 LLM 大部分计算花在大型矩阵乘法(GEMM)和相关注意力内核。通过提供高度优化、Ascend 原生且镜像广泛使用的 DeepGEMM API 的实现,此项目实现:
- 无缝移植现有基于 PyTorch 的模型到 Ascend,无需重写内核。
- 准确的性能基准,供研究人员比较 Ascend 与 NVIDIA/AMD GPU。
- 参考代码,供硬件工程师探索新的 Ascend 优化技术(稀疏加载、协程流水线、自定义缩放因子打包)。
许可证与引用
- 许可证: MIT – 允许商业和研究使用。
- 引用: 提供可直接复制的 BibTeX 条目供学术论文使用。
TL;DR
DeepGEMM‑Ascend 是可直接替换、高性能的华为 Ascend NPU GEMM 库,暴露与原始 DeepGEMM 项目相同的 Python API,同时加入 Ascend 特定优化。它针对 LLM 工作负载(密集 GEMM、MoE、MQA),并声称在各种矩阵形状下接近峰值硬件利用率。可通过 pip 安装,需要 Ascend CANN 工具包和 C++20 编译器。
相关
- 项目
- 项目
- 项目
- 项目