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」内核的详细数据。

如何开始(快速入门摘要)

  1. 硬件与软件先决条件
    • 华为 Ascend NPU(已在 950 系列验证)
    • CANN 9.20 工具包(bisheng、ld.lld)
    • torch_npu Python 包
    • Python 3.10+、兼容 C++20 的编译器,以及 tilelang、tree‑sitter、tree‑sitter‑cpp
  2. 使用子模块克隆
    git clone --recursive https://github.com/deepseek-ai/DeepGEMM-Ascend.git
    cd DeepGEMM-Ascend
    
  3. 构建 C++ 扩展(仓库提供 develop.sh 脚本,会链接 include 并运行编译器)。
  4. 安装
    pip install . --no-build-isolation
    
  5. 在 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 编译器。

相关

  • 项目
  • 项目
  • 项目
  • 项目