deepseek-ai/DeepGEMM-Ascend

What is DeepGEMM‑Ascend?

DeepGEMM‑Ascend is a hardware‑specific port of the open‑source DeepGEMM matrix‑multiply library, targeting Huawei’s Ascend NPU family (e.g., Ascend 950). It implements the same Python‑level API as DeepGEMM, so code written for other platforms (CUDA, ROCm, etc.) can run on Ascend simply by installing this package.

The library wraps the Ascend MAD (matrix‑add‑multiply) primitives with a thin C++/Python extension, handling the low‑level details of fractal layouts, alignment, and address calculations. It also adds Ascend‑only optimisations such as sparse data loading and coroutine‑based pipelining, aiming to reach the hardware’s theoretical peak performance for a variety of matrix shapes.

---

Who might use it?

Audience	Why it matters
AI researchers / engineers building large language models (LLMs) that run on Ascend NPUs – they need fast GEMM, MQA, and MoE kernels.
Framework developers extending PyTorch‑NPU (torch_npu) with custom high‑performance ops.
Performance engineers looking for reference implementations of extreme‑optimised Ascend kernels.
Hardware‑aware ML practitioners who want to benchmark BF16, FP8, FP4, and mixed‑precision matrix ops on Ascend.

---

Core features (as described in the README)

Feature	Detail
API compatibility	Uses the same package name and function signatures as the upstream DeepGEMM project.
Precision support	BF16, FP8, FP4 GEMM; also specialised kernels for MQA logits and MegaMoE.
Lightweight abstraction	Hides Ascend‑specific layout and alignment complexities while keeping the kernel code concise.
Ascend‑specific optimisations	Sparse data loading, coroutine‑based pipelining, custom scaling‑factor packing (UE8M0 pairs → int16).
Utility helpers	Functions to set number of AI cores, choose NPU architecture, control deterministic mode, transform scaling‑factor layouts, and query theoretical alignment limits.
JIT runtime	Relies on the DeepJIT runtime; kernels are compiled at first use and cached (configurable via DG_JIT_* env vars).
Environment‑variable tuning	Debug, cache location, launch timeout, and other knobs exposed via DG_JIT_* and standard Ascend variables.
Performance	Benchmarks show 98‑99 % of the hardware limit for dense GEMM across BF16/FP8/FP4, plus detailed numbers for grouped GEMM, MQA, MegaMoE, and a custom “HC Prenorm” kernel.

---

How to get started (quick‑start summary)

1. Hardware & software prerequisites
   • Huawei Ascend NPU (validated on 950 series)
   • CANN 9.20 toolkit (bisheng, ld.lld)
   • torch_npu Python package
   • Python 3.10+, a C++20‑compatible compiler, plus tilelang, tree‑sitter, tree‑sitter‑cpp
2. Clone with submodules

   git clone --recursive https://github.com/deepseek-ai/DeepGEMM-Ascend.git
   cd DeepGEMM-Ascend
   

3. Build the C++ extension (the repo provides a develop.sh script that links includes and runs the compiler).
4. Install

   pip install . --no-build-isolation
   

5. Use the same DeepGEMM API in your Python code, e.g.

   import deep_gemm as dg
   dg.set_num_sms(8)          # choose how many AI cores to use
   out = dg.aclnn_fp8_fp4_gemm_nt(A, B, scaling_factors)
   

---

When does it shine?

The README lists several benchmark categories where the library approaches the Ascend hardware ceiling:

• Dense GEMM – up to 99.8 % of the theoretical TFLOPS for BF16, FP8, FP4.
• M‑Grouped GEMM – useful for Mixture‑of‑Experts (MoE) experts, achieving > 800 TFLOPS on realistic expert‑token patterns.
• MQA Logits – fast scoring for multi‑query attention, with latency in the 100‑300 µs range.
• MegaMoE – end‑to‑end fused kernels (dispatch + two GEMMs + SwiGLU) that keep both compute and memory bandwidth high.
• HC Prenorm GEMM – a specialised kernel for the “Manifold‑Constrained Hyper‑Connections” module, saturating HBM write bandwidth.

These numbers suggest the library is ready for production‑scale LLM inference/training workloads on Ascend hardware.

---

Limitations / things to be aware of

• Hardware‑specific – only works on Huawei Ascend NPUs; no CPU or GPU fallback.
• Toolchain‑heavy – requires CANN 9.20, bisheng compiler, and a C++20 compiler with <format> support.
• Environment‑variable tuning – performance may depend on correctly setting DG_JIT_* variables and the Ascend launch‑blocking/debug flags.
• Beta status – the repository reports an initial release (2026‑09‑30); community support and documentation may still be limited.

---

How it fits into the AI ecosystem

DeepGEMM‑Ascend sits at the low‑level linear‑algebra layer of the AI stack. Modern LLMs spend the majority of their compute in large matrix multiplications (GEMM) and related attention kernels. By providing a highly optimised, Ascend‑native implementation that mirrors the widely‑used DeepGEMM API, the project enables:

• Seamless porting of existing PyTorch‑based models to Ascend without rewriting kernels.
• Accurate performance baselines for researchers comparing Ascend against NVIDIA/AMD GPUs.
• Reference code for hardware engineers exploring new Ascend optimisation techniques (sparse loading, coroutine pipelines, custom scaling‑factor packing).

---

License & citation

• License: MIT – permissive for commercial and research use.
• Citation: A ready‑to‑copy BibTeX entry is provided for academic papers.

---

TL;DR

DeepGEMM‑Ascend is a drop‑in, high‑performance GEMM library for Huawei Ascend NPUs, exposing the same Python API as the original DeepGEMM project while adding Ascend‑specific optimisations. It targets LLM workloads (dense GEMM, MoE, MQA) and claims near‑peak hardware utilisation across a wide range of matrix shapes. Installable via pip, it requires the Ascend CANN toolkit and a C++20 compiler.