DeepGEMM Ascend: DeepGEMM kernels rebuilt for Huawei Ascend NPUs

What DeepGEMM Ascend is and who needs it

DeepGEMM Ascend is a port of the DeepGEMM matrix multiplication kernel library to the Huawei Ascend platform. The README states that it is fully API-compatible with DeepGEMM and that the two projects use the same package name, so a user installs the package and keeps the same APIs and development workflow used on other supported platforms. That compatibility claim is the whole reason the project exists, and it is also the reason to read the rest of this article carefully: what is portable is the interface, not every numerical detail.

The feature list covers BF16, FP8 and FP4 GEMM, MQA logits and MegaMoE operators. The target audience is narrow: teams deploying DeepSeek-style models on Ascend NPUs who want the tuned kernels rather than a generic PyTorch path, and who would rather not rewrite their kernel calls when they move hardware.

The project is young. The News section in the README records a single entry, the initial release on 2026-09-30, with support for Ascend 950 devices and kernels designed to achieve near-peak hardware performance. The repository was last pushed on 2026-10-14 and is not archived, and it carries an MIT licence, which is a permissive choice for a library that other vendors' code may depend on.

The structure is small and conventional for a kernel project: a csrc directory for the C++ sources, a deep_gemm directory for the Python package, a scripts directory, a tests directory, a third-party directory for vendored dependencies, plus setup.py, build.sh, develop.sh, CMakeLists.txt and a .clang-tidy configuration. The third-party directory is where the JIT runtime and the enum library live, since the build include paths in setup.py point at third-party/deep-jit and third-party/magic_enum.

The Ascend abstraction: MAD primitives, fractal layouts and alignment

The interesting engineering decision in this port is the abstraction layer. DeepGEMM Ascend provides a lightweight abstraction over the Ascend MAD, or matrix multiply-add, primitives. According to the README, this hides much of the complexity associated with fractal layouts, alignment constraints, address calculations and verbose low-level parameters, and the result is that GEMM kernels stay both concise and efficient.

The four hidden categories are the ones that usually dominate a hand-written Ascend kernel. Fractal layout means the data layout the hardware expects is not the row-major layout a Python caller assumes, so address arithmetic has to be rewritten. Alignment constraints mean tensor dimensions and strides must satisfy hardware rules that are not visible in the API. The low-level parameters are the part where a naive port accumulates dozens of positional arguments that the abstraction collapses.

On top of that abstraction the port applies Ascend-specific techniques, and the README names two: sparse data loading and coroutine-based pipelining. The intent is to approach the performance limits of the Ascend hardware rather than to sit at a comfortable distance from them. The README also makes a claim about the implementation style, that the codebase is lightweight and reaches hardware-limit performance across a wide range of matrix shapes, and it adds that these implementations can serve as references for extreme performance optimization on the Ascend platform. That is a project positioning itself as a teaching example as much as a library.

Building the extension and installing the package

Development starts with a recursive clone, because the README notes that the submodule must be cloned.

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

The develop script is the build entry point for development, and the README's instruction to look at it before running it is worth heeding, since it is what links the includes and builds the C++ extension. Installation of the built package then uses pip with build isolation turned off.

pip install . --no-build-isolation

The requirements list explains why the build is not a plain pip install. You need an HUAWEI Ascend NPU, developed and validated on the Ascend 950 series. The CANN 9.20 toolkit has to provide bin/bisheng and bin/ld.lld. The torch_npu package is required. Python 3.10 or higher is needed, along with compilers and standard libraries with C++20 `<format>` support. tilelang is used by the HC prenorm kernel and is declared as a package dependency, and tree-sitter with tree-sitter-cpp are used to generate Python type stubs when building from source.

The setup.py file shows the rest of the plumbing. It reads the toolkit location from ASCEND_HOME_PATH with a fallback to ASCEND_TOOLKIT_HOME and then to /usr/local/Ascend/ascend-toolkit/latest. The compiler flags are `-std=c++20`, `-O3`, `-g1` and `-fPIC`, and the extension links ascendcl, hccl_fwk, torch_npu, opapi and dl. Two build switches exist, DG_SKIP_BUILD and DG_USE_LOCAL_VERSION, the latter defaulting to 1.

Kernel interfaces and the scaling factor layout that differs from NVIDIA

The kernel API follows DeepGEMM, and the README points readers at the upstream DeepGEMM interface documentation rather than duplicating it. Two interface families are named directly.

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

The `bf16_gemm_nn` and `fp8_fp4_gemm_nn` interfaces take A with shape M by K and B with shape K by N. For dense BF16, FP8, FP4 and mixed FP8 by FP4 GEMMs, the output tensor c may have shape M by N or 1 by N, with the same dtype as d. A full C tensor is accumulated with atomic add, and a contiguous 1 by N bias is loaded and added in the AIV epilogue before the output cast.

The one place where the port is not a drop-in replacement is the scaling factor format, and the README flags it as a note. On Ascend, 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 hardware efficiency. That is a different memory layout from the NVIDIA path, so code that manipulates scaling factors directly, or code that assumes the reference layout, has to be adjusted. The utility functions exist for exactly this reason, with transform_sf_into_required_layout and transform_k_grouped_sf_into_required_layout for the two cases.

Utility functions and the JIT environment variables

Beyond the kernels, the package exposes a set of controls that tune how the runtime behaves on the device. The number of AI cores the kernels may use is set with set_num_sms and read back with get_num_sms. The dav-architecture string used for JIT compilation can be overridden with set_npu_arch and queried with get_npu_arch, and passing 0 queries the device for it. Contiguous layout alignment has a setter, a getter and a separate query for the theoretical minimum M and K alignment. Deterministic algorithms can be switched on and off. The paged MQA kernels have their scheduling metadata built by get_paged_mqa_logits_metadata, and a family of aclnn reference GEMMs, aclnn_fp8_fp4_gemm_{nt, nn, tn, tt} and aclnn_bf16_gemm_{nt, nn, tn, tt}, exist to cross-check the kernels in tests.

The environment variables control the JIT runtime rather than the kernels themselves, and the fallback rule is worth noting: every DG_JIT_ variable falls back to the corresponding global DJ_JIT_ variable when unset, and the values are snapshotted when the JIT runtime is first created. Changing a variable after the runtime exists therefore has no effect.

The cache behaviour is the one that will affect daily use. DG_JIT_CACHE_DIR takes a directory or a colon-separated list of directories, the lookup searches all paths front to back with the first hit winning, and a cache miss compiles into the first path, with $HOME/.dj as the default. The remaining variables are diagnostic switches. DG_JIT_DEBUG turns on compiler commands, load-time reporting and the selected config per shape, and DG_PRINT_CONFIGS prints the selected config for each shape. DG_JIT_PRINT_COMPILER_COMMAND, DG_JIT_PRINT_LOAD_TIME and DG_JIT_KERNEL_DEBUG_INFO cover compiler and disassembler commands, kernel load time and Bisheng kernel debug information. DG_JIT_LAUNCH_TIMEOUT is a launch timeout in seconds that defaults to 300 and disables at 0, and ASCEND_LAUNCH_BLOCKING synchronizes the device after every launch, turning asynchronous launch errors into in-line failures.

Measured performance on Ascend 950 and what the tables leave out

The performance section of the README is specific about its conditions. Measurements were taken on Ascend 950DT hardware with CANN 9.20 using bench_msprof with a cold L2, and the shapes and cases follow the DeepGEMM test suite, covering inference and training workloads of the DeepSeek model series. The tests directory is named as the place to reproduce them.

For dense GEMM the reported figures are given as a share of the hardware limit, and the highest is BF16 by BF16 at 99.8 percent: a shape of 4096 by 7168 by 16384 at 2229.9 microseconds, 431 TFLOPS against a 432 TFLOPS limit. FP8 by FP8 and FP8 by FP4 both sit at 99.5 percent, 861 TFLOPS against 865, at just over 1117 microseconds. FP4 by FP4 is listed at 98.3 percent, 1701 TFLOPS against a 1730 limit, in 565.6 microseconds. A second table gives inference-oriented FP8 shapes with latency, computation and memory bandwidth figures, including a 128 by 32768 by 512 case at 10.4 microseconds and 1829 GB/s, and a 128 by 7168 by 16384 case at 54.8 microseconds and 2456 GB/s.

What the tables do not include matters as much as what they do. There is no comparison against a baseline such as a vendor library or an unfused PyTorch path, no figures for the MQA logits or MegaMoE operators, and no data on hardware other than Ascend 950. The claimed advantage, reaching the hardware limit, is a claim about the hardware ceiling rather than about the software being faster than an alternative, which is a much harder bar to clear on a shared cluster. If your workload is dominated by small shapes, by attention rather than GEMM, or by a generation device other than 950DT, these numbers say little about your case.

Editorial conclusion

DeepGEMM Ascend is the right choice for a team that already runs DeepGEMM code on other platforms and now needs the same API on Ascend hardware, with a CANN 9.20 toolkit and a torch_npu stack in place. It is the wrong choice for a project without Ascend hardware, since the kernels are built and JIT-compiled for the target device. Verify first that your NPU is on the Ascend 950 series the project was developed against, that the CANN toolkit provides bin/bisheng and bin/ld.lld, and that your scaling factor handling accounts for the UE8M0 packing rule that differs from the NVIDIA path.

Frequently asked questions

Official sources

1. deepseek-ai/DeepGEMM-Ascend on GitHub
2. Issues
3. License: MIT
4. README