DeepGEMM-Ascend README

The project's README from github.com/deepseek-ai/DeepGEMM-Ascend. As one markdown file: README.md.

DeepGEMM Ascend

DeepGEMM Ascend is a port of DeepGEMM to the HUAWEI Ascend platform. It is fully API-compatible with DeepGEMM and supports BF16, FP8, FP4 GEMM, MQA logits, and MegaMoE. On Ascend platforms, users can simply install the package and use the same APIs and development workflow as DeepGEMM on other supported platforms.

DeepGEMM Ascend provides a lightweight abstraction over the Ascend MAD (matrix multiply-add) primitives, hiding much of the complexity associated with fractal layouts, alignment constraints, address calculations, and verbose low-level parameters. This enables GEMM kernels to remain both concise and efficient. DeepGEMM Ascend makes extensive use of Ascend-specific optimization techniques, such as sparse data loading and coroutine-based pipelining, to approach the performance limits of the Ascend hardware. These implementations can also serve as references for extreme performance optimization on the Ascend platform.

Despite its lightweight codebase, DeepGEMM Ascend can achieve peak hardware performance across a wide range of matrix shapes.

DeepGEMM Ascend 是华为昇腾平台上的 DeepGEMM 实现。它完全兼容 DeepGEMM 的 API，支持 BF16、FP8、FP4 GEMM、MQA logits 和 MegaMoE 算子。DeepGEMM Ascend 和 DeepGEMM 使用相同的包名，用户只需安装当前包，即可沿用其他平台上 DeepGEMM 的 API 和开发流程。

DeepGEMM Ascend 对昇腾平台的矩阵乘加原语（MAD）提供了一个轻量的抽象，能够隐藏分形矩阵布局、对齐约束、地址计算、参数转换等细节，让 GEMM kernel 的实现保持简洁高效。DeepGEMM Ascend 广泛采用昇腾平台特有的优化技术，包括稀疏数据加载、基于协程的流水线等，以接近硬件的性能极限。这些实现也可以作为昇腾平台极致性能优化的参考。

DeepGEMM Ascend 代码轻量，并且在多种矩阵形状上能达到硬件极限性能。

News

Quick Start

Requirements

Development

# Submodule must be cloned
git clone --recursive https://github.com/deepseek-ai/DeepGEMM-Ascend.git
cd DeepGEMM-Ascend

# Link some essential includes and build the C++ extension
cat develop.sh
./develop.sh

Installation

pip install . --no-build-isolation

Interfaces

Kernel Interface

Please refer to DeepGEMM's interfaces for the kernel APIs.

[!NOTE] The scaling factor format on Ascend differs from NVIDIA's: each pair of UE8M0 scaling factors along the K dimension is packed into an int16, and the packed values are stored in MN-major order for optimal hardware efficiency.

Utilities

The library provides some utility functions besides the above kernels:

Environment Variables

Ascend home path is given by ASCEND_HOME_PATH or ASCEND_TOOLKIT_HOME.

Each DG_JIT_* variable falls back to the corresponding global DJ_JIT_* variable when unset, and is snapshotted when the JIT runtime is first created.

Performance

Measured on Ascend 950DT (CANN 9.20) using bench_msprof with cold L2. Shapes and cases follow the DeepGEMM test suite and cover inference and training workloads of the DeepSeek model series; see tests/ for reproduction.

Dense GEMM

Dense GEMM reaches up to 99.8% of the hardware limit across types:

Type M N K Latency (us) Computation (TFLOPS) Hardware limit (TFLOPS) Utilization (%)
FP4×FP4 4096 7168 16384 565.6 1701 1730 98.3
FP8×FP4 4096 7168 16384 1117.8 861 865 99.5
FP8×FP8 4096 7168 16384 1117.6 861 865 99.5
BF16×BF16 4096 7168 16384 2229.9 431 432 99.8

FP8 GEMM for inference:

M N K Latency (us) Computation (TFLOPS) Memory bandwidth (GB/s)
128 2112 7168 15.0 258 1205
128 576 7168 10.1 105 564
128 24576 1536 20.6 468 2068
128 32768 512 10.4 415 1829
128 7168 16384 54.8 549 2456
128 4096 7168 19.0 397 1797
128 7168 2048 10.1 373 1670
4096 2112 7168 156.9 790 319
4096 576 7168 54.6 619 689
4096 24576 1536 362.2 854 137
4096 32768 512 164.3 837 129
4096 7168 16384 1117.6 861 186
4096 4096 7168 282.9 850 234
4096 7168 2048 143.4 838 181

M-Grouped GEMM

Grouped GEMM for DeepSeek MoE experts (FP8×FP4, BF16 output). #Groups is the number of experts, and M per group the average tokens each receives.

#Groups M per group N K Latency (us) Computation (TFLOPS) Memory bandwidth (GB/s)
4 8192 6144 7168 3722.4 818 104
4 8192 7168 3072 1778.0 856 98
4 8192 4096 4096 1356.1 855 148
4 8192 4096 2048 680.6 852 148
8 4096 6144 7168 3689.4 831 136
8 4096 7168 3072 1778.9 862 130
8 4096 4096 4096 1356.5 861 180
8 4096 4096 2048 681.0 858 179

MQA Logits

MQA scoring for the DeepSeek Lightning Indexer. The kernel is FIX-pipe bound rather than compute bound, saturating the FIX pipe at 99% utilization.

Type Format #Q Tokens #K Tokens #Heads Dim Latency (us) Computation (TFLOPS) Memory bandwidth (GB/s)
Prefill FP8 4096 8192 32 128 302.6 681 283
Prefill FP4 4096 8192 32 128 277.4 743 277
Decode FP8 256 8192 32 128 150.9 565 2006
Decode FP4 256 8192 32 128 124.2 708 1349

MegaMoE

Mega MoE fuses EP dispatch, two grouped GEMMs, SwiGLU, and combine. Benchmarked over EP8 with top-k=6 (each token routed to 6 experts) and one shared expert; all values are averaged across 8 ranks.

#Experts Hidden Intermediate Tokens Latency (us) Computation (TFLOPS) Memory bandwidth (GB/s) Communication bandwidth (GB/s)
96 5120 2304 64 138.7 228.5 1850.0 37.5
96 5120 2304 256 225.5 562.5 1257.4 91.6
96 7168 3072 64 222.2 266.4 2136.9 32.8
96 7168 3072 256 359.2 659.2 1425.4 80.5
384 5120 2304 4096 2567.7 790.3 567.5 128.8
384 5120 2304 16384 9768.6 831.0 325.0 135.2
384 7168 3072 4096 4600.7 823.4 531.3 100.6
384 7168 3072 16384 17904.2 846.3 269.3 103.3

HC Prenorm GEMM

Prenorm GEMM for the DeepSeek mHC (Manifold-Constrained Hyper-Connections) module, which nearly saturates the HBM write bandwidth.

M N K Latency (us) Computation (TFLOPS) Memory bandwidth (GB/s)
13 24 28672 6.7 3 519
137 24 28672 9.6 20 1110
512 24 28672 15.4 46 2082
4096 24 28672 72.7 78 3276
8192 24 28672 136.7 82 3463

Contributors

Acknowledgements

DeepGEMM-Ascend follows the design of the upstream DeepGEMM project, which is inspired by CUTLASS. The JIT runtime is provided by DeepJIT. The mHC kernel is backed by Tilelang. We sincerely thank the developers of these projects for their contributions, and gratefully acknowledge Huawei for its technical support and engineering expertise throughout the development of DeepGEMM-Ascend.

License

This code repository is released under the MIT License.

Citation

If you use DeepGEMM-Ascend in your work, please cite:

@misc{deepgemm_ascend2026,
  title     = {DeepGEMM-Ascend: Clean and Efficient BLAS Kernel Library on Ascend NPU},
  author    = {Kexing Zhou and Zhean Xu and Anyi Xu and Chenggang Zhao and
               Yuxuan Zhou and Yunfan Xiao and Guanglin Li and Kaifeng Chen and Yuhao Meng and
               Huanqi Cao and Ruifan Xu and Chenhao Xu and Kuai Yu},
  year      = {2026},
  publisher = {GitHub},
  url       = {https://github.com/deepseek-ai/DeepGEMM-Ascend}
}

DeepGEMM-Ascend (deepseek-ai/DeepGEMM-Ascend)

Searching Database…

583 ·GitHub
DeepGEMM-Ascend
这是什么？DeepGEMM-Ascend 是把 DeepSeek 的 DeepGEMM 矩阵乘法内核库移植到华为昇腾 NPU 的项目，与原版完全 API 兼容，支持 BF16、FP8 和 FP4 GEMM，以及 MQA logits 和 MegaMoE 算子。
这是一个独立应用，还是我引入代码里的库？它是一个底层内核库（没有 UI 或 CLI），对昇腾的矩阵乘加（MAD）原语提供轻量抽象，隐藏分形布局、对齐约束和地址计算，让 GEMM 内核保持简洁和高速。
怎么让它跑起来？在配好昇腾工具链的机器上安装这个包（名称和 API 与上游 DeepGEMM 相同），然后调用相同的内核函数，遵循与其他 DeepGEMM 支持平台一致的开发流程。
试用之前需要什么硬件和软件？需要一块华为昇腾 NPU（在昇腾 950 系列上开发和验证）、CANN 9.20 工具包（bisheng 编译器、ld.lld）、torchnpu 包、Python 3.10+，以及支持 C++20 的编译器；HC prenorm 内核需要 tilelang，从源码构建 Python 类型存根时需要 tree-sitter/tree-sitter-cpp。
它在 NVIDIA GPU 上也能用吗？不能，它专为昇腾硬件构建，无法在 NVIDIA GPU 上运行；其设备上的缩放因子布局甚至与 NVIDIA 的不同，因此针对一个平台调优的代码需要用对应的 DeepGEMM 变体。
它是为谁打造的？在华为昇腾硬件上运行 DeepSeek 类模型推理或训练负载（稠密 GEMM、MoE 分组 GEMM、MQA 注意力、mHC）、想要 DeepGEMM 原样 API 而不必为 NVIDIA 重写内核的工程师。
MegaMoE 是什么？MegaMoE 是一个融合内核，把专家并行（EP）分发、两个分组 GEMM、SwiGLU 激活和 combine 步骤合并用于混合专家推理，基准测试采用 EP8、top-k=6 路由加一个共享专家。
它能达到硬件峰值性能的多大比例？在昇腾 950DT 上，稠密 GEMM 在 FP4、FP8 和 BF16 数据类型下最高达到理论硬件极限的 99.8%。
为什么要做这个项目，而不直接用 DeepGEMM？上游 DeepGEMM 面向 NVIDIA GPU；这个移植版在华为昇腾的 MAD 原语上重新实现了同样的内核，让现有基于 DeepGEMM 的代码无需改动就能在昇腾硬件上运行。
HC Prenorm GEMM 内核是做什么用的？它实现 DeepSeek mHC（Manifold-Constrained Hyper-Connections）模块所用的 prenorm GEMM，设计目标是几乎打满昇腾的 HBM 写带宽。
What are similar projects?DeepEP-Ascend, DeepGEMM, TensorRT-LLM, tilelang, DeepEP, FlashMLA, DeepSeek-V3, cutlass
Which topics is it in?cpp20, linear-algebra, jit, mixture-of-experts, moe, llms, kernel, optimization, low-level, llm-inference, high-performance, compiler
Creator
deepseek-ai
Size
0.3 MB
Stars
583
Updated
CommitOct 14
Age
18d
Commits · 90 Days