DeepGEMM
DeepGEMM is an open-source tensor-core kernel library from DeepSeek. It began as a set of FP8 general matrix multiplication (GEMM) kernels for NVIDIA Hopper GPUs, released on 26 February 2025 as the third project in DeepSeek's "Open Source Week," a five-day series starting 24 February 2025 in which the company published several pieces of the training and inference infrastructure behind DeepSeek-V3 and DeepSeek-R1 [1][2][4][9]. DeepSeek introduced it as "an FP8 GEMM library that supports both dense and MoE GEMMs, powering V3/R1 training and inference," reaching "Up to 1350+ FP8 TFLOPS on Hopper GPUs" with "Core logic at ~300 lines" in code the company called "as clean as a tutorial" [2][15].
Since then the project has grown well past a single kernel. As of 1 October 2026 the repository describes DeepGEMM as "a unified, high-performance tensor core kernel library that brings together the key computation primitives of modern large language models," covering GEMMs in FP8, FP4 and BF16, a fused mixture-of-experts kernel with overlapped communication called Mega MoE, multi-query-attention scoring for the sparse "lightning indexer" of DeepSeek's later models, and HyperConnection kernels, all in one CUDA codebase compiled at runtime [1]. The GitHub repository carried 7,906 stars and 1,285 forks on 1 October 2026, and the installed Python package reports version 2.8.1 [18][28]. On 30 September 2026 DeepSeek published DeepGEMM-Ascend, an API-compatible port to Huawei Ascend NPUs, as part of a coordinated release of its infrastructure stack for Huawei hardware [19][22].
DeepGEMM matters because matrix multiplication dominates the cost of both training and serving large language models, and doing that arithmetic in 8-bit floating point (FP8) roughly doubles throughput and halves memory traffic compared with 16-bit formats [8]. The hard part is keeping FP8 numerically stable. DeepGEMM packages the techniques DeepSeek used to train a frontier model in FP8 into a clean, reusable component [1][3].
Is DeepGEMM open source?
Yes. DeepGEMM is released under the permissive MIT license, and its full source lives in DeepSeek's public GitHub repository [1][28]. Installation compiles only a host-side extension, never the kernels: all kernels are compiled at runtime, so the library only ever compiles the exact matrix shapes a workload needs [1]. Since the September 2026 release that compilation goes through DeepJIT, a separate DeepSeek repository published on 8 September 2026 and described as "A lightweight library for xPU kernel JIT compilation," also MIT-licensed [1][29].
At launch DeepSeek presented the library as a teaching reference as much as a production tool: the February 2025 README said it was "designed for simplicity, with only one core kernel function comprising around ~300 lines of code," which made it "a clean and accessible resource for learning Hopper FP8 matrix multiplication and optimization techniques" [16]. The current README keeps the framing but drops the line count, saying the library has "only a limited number of core kernel functions" and is a resource "for learning NVIDIA GPU kernel optimization techniques" [1].
Background and release
DeepSeek-V3 was one of the first widely deployed large models trained natively in FP8. Its technical report describes a fine-grained quantization scheme in which activations and weights are scaled in small blocks or tiles rather than per tensor, so that a few large values do not force the rest of a tensor into a lossy range [3]. Activations are grouped and scaled on a 1x128 tile basis (one token across 128 channels) and weights on a 128x128 block basis (128 input channels by 128 output channels), with the E4M3 FP8 format used on all tensors [3]. Reproducing that scheme efficiently requires GEMM kernels that understand per-block scaling factors and that compensate for the limited accumulation precision of Hopper's FP8 Tensor Cores.
The full V3 model is large: 671 billion total parameters with 37 billion activated per token, pre-trained on 14.8 trillion tokens [3]. DeepSeek reported the pre-training at 2.664 million H800 GPU hours and full training at 2.788 million H800 GPU hours, unusually low figures for a model of that scale [3]. Efficient FP8 matrix multiplication is one reason those numbers are as small as they are, since GEMM is where most of the floating-point work happens.
Rather than keep those kernels internal, DeepSeek open-sourced them during Open Source Week, a five-day run that began on 24 February 2025. The series included FlashMLA on day one (an attention-decoding kernel for Hopper), DeepEP on day two, which DeepSeek introduced as "the first open-source EP communication library for MoE model training and inference," DeepGEMM on day three, the DualPipe pipeline-parallelism algorithm and the EPLB expert-parallel load balancer on day four, and the Fire-Flyer File System (3FS) on day five [2][10]. The open-infra-index that catalogs the week describes DeepGEMM as "an FP8 GEMM library that supports both dense and MoE GEMMs, powering V3/R1 training and inference," with "Up to 1350+ FP8 TFLOPS on Hopper GPUs" [2]. The GitHub repository was created on 13 February 2025 and its first public commit landed on 25 February 2025, the day before the announcement [28].
How does DeepGEMM keep FP8 accurate?
FP8 stores each number in eight bits, which gives very little room for both range and precision. DeepGEMM's launch version used the E4M3 variant (four exponent bits, three mantissa bits) for the matrix inputs, matching DeepSeek-V3's training recipe [3][8][12][16]. Multiplying two FP8 matrices is fast on Hopper Tensor Cores, but two numerical problems get in the way. The first is dynamic range: a single tensor often contains values spanning many orders of magnitude, and one global scale cannot represent all of them well. DeepGEMM addresses this by consuming the fine-grained, per-tile and per-block scaling factors that DeepSeek's quantization recipe produces, so each small region of a matrix carries its own scale [1][3].
The second problem is accumulation precision. A GEMM sums many products into each output element, and the V3 technical report observes that "the accumulation precision of FP8 GEMM on NVIDIA H800 GPUs is limited to retaining around 14 bits, which is significantly lower than FP32 accumulation precision," enough to produce a maximum relative error of nearly 2 percent in a test with an inner dimension of 4,096 [3]. DeepSeek's later hardware paper spells out the mechanism: after aligning the mantissa products, the Tensor Core keeps only their highest 13 fraction bits and accumulates into FP22 registers with one sign bit, eight exponent bits and 13 mantissa bits [13]. This is a property of the hardware, not a choice DeepSeek made.
DeepGEMM's answer is two-level accumulation, also called promotion. The Tensor Cores do the bulk of the multiply-accumulate work quickly, and once an interval of 128 elements (four WGMMA instructions) is reached, the partial results are copied to FP32 registers on the CUDA cores, where the running sum is kept at full precision [1][3][16]. DeepSeek reports 128 elements as "the minimal accumulation interval that can significantly improve precision without introducing substantial overhead" [3]. The result is FP8 throughput with numerical behavior close enough to higher precision to train and serve a frontier model.
The scaling-factor contract has become more formal over time. A design document in the repository specifies that every FP32 scaling factor must be an exact power of two, that SM90 kernels take scaling factors in FP32 while SM100 kernels take them packed as UE8M0 exponents (four per 32-bit integer), and that the storage granularity along K may be 32 or 128 on SM100 but only 128 on SM90 [1][17].
How is DeepGEMM optimized for Hopper GPUs?
DeepGEMM was written for the NVIDIA Hopper architecture and leans on hardware features specific to it. The most important is the Tensor Memory Accelerator (TMA), a dedicated unit introduced with Hopper that moves large tiles of data between global and shared memory asynchronously, without tying up the regular compute threads [6][14]. The launch README lists TMA loads for the left-hand matrix, its scaling factors and the right-hand matrix, TMA stores for the output, TMA multicast for the left-hand matrix, and TMA descriptor prefetching [16]. The scaling-factor layouts the kernels expect are still described as "TMA-aligned" for the same reason [1].
On top of the TMA, the library uses a persistent, warp-specialized kernel design. Different warps are assigned to distinct jobs: some handle the TMA data movement, some issue the matrix-multiply-accumulate (MMA) instructions on the Tensor Cores, and some run the CUDA-core promotion step. Keeping these stages overlapping means the memory pipeline and the compute pipeline run at the same time rather than waiting on each other [16]. Two smaller tricks appeared in the launch write-up. DeepGEMM supported unaligned block sizes (for example a block of 112 rather than a power of two), which improves streaming-multiprocessor utilization on matrix shapes that do not divide evenly into standard tiles [16]. It also interleaved FFMA (fused floating-point multiply-add) instructions at the SASS assembly level, a low-level scheduling tweak DeepSeek reported as "10%+ in some cases" [16]. That post-compilation step was retired in the July 2025 refactor, because, as the README puts it, "NVCC 12.9 will automatically do the FFMA interleaving" [1].
Just-in-time compilation is itself a design pillar. All kernels are compiled at runtime, so there is no CUDA build step at installation and the library only ever compiles the shapes a workload needs [1]. An update on 7 May 2025 added an NVRTC path that DeepSeek reported as "up to 10x compilation speedup," enabled with the DG_JIT_USE_NVRTC environment variable [32]. The July 2025 refactor disabled that path in favour of a low-CPU-overhead JIT C++ module, and by the September 2026 release the JIT runtime had been split out into DeepJIT [1][29].
How does DeepGEMM support mixture-of-experts models?
Mixture-of-experts models do not run one big matrix multiply per layer. Instead each token is routed to a subset of experts, and every expert multiplies only the tokens sent to it, so a layer turns into many smaller GEMMs of varying sizes. DeepGEMM provides grouped GEMM kernels for exactly this pattern, in two layouts [1].
The contiguous layout is aimed at training forward passes and inference prefilling. Tokens for the different experts are concatenated along the M axis (the row dimension) while the N and K dimensions stay fixed, and each expert segment must be aligned to the GEMM M block size [1]. Unlike CUTLASS grouped GEMMs, DeepGEMM groups only the M axis, a design the README says is "tailored for scenarios where experts in an MoE model share the same shape" [1]. The masked layout is aimed at decoding-time inference. The number of tokens per expert is not known until routing happens at runtime, which is awkward when the kernel launch is captured ahead of time in a CUDA graph. DeepGEMM handles this with a mask tensor that tells the kernel which portions are valid, so it computes only the real work while still fitting inside a static CUDA graph, and the README suggests feeding it the output of DeepEP's low-latency kernels [1]. A K-axis-grouped API was later added for MoE weight backward passes, with M and N fixed [1].
How fast is DeepGEMM?
DeepSeek benchmarked the launch version on NVIDIA H800 SXM5 GPUs with NVCC 12.8, testing "all shapes potentially used in DeepSeek-V3/R1 inference (including both prefilling and decoding, but without tensor parallelism)," and compared the results against "our internally and carefully optimized implementation based on CUTLASS 3.6" [16]. The highest dense figure in that table was 1,358 TFLOPS on a 4096x7168x16384 multiplication; the highest memory bandwidth, 2,668 GB/s, came from a much smaller 64x7168x16384 shape that is bandwidth-bound rather than compute-bound [16]. Dense speedups over the baseline ran from 1.0x to 2.7x, with the largest gains on the small, irregular shapes where generic libraries leave performance unused, and the grouped MoE kernels in both layouts ran 1.1x to 1.2x faster than the baseline [16].
| M x N x K | Kernel type | Computation | Memory bandwidth | Speedup |
|---|---|---|---|---|
| 64 x 2112 x 7168 | Dense GEMM | 206 TFLOPS | 1,688 GB/s | 2.7x |
| 64 x 7168 x 16384 | Dense GEMM | 336 TFLOPS | 2,668 GB/s | 1.4x |
| 128 x 7168 x 16384 | Dense GEMM | 645 TFLOPS | 2,604 GB/s | 1.4x |
| 4096 x 4096 x 7168 | Dense GEMM | 1,304 TFLOPS | 500 GB/s | 1.1x |
| 4096 x 7168 x 16384 | Dense GEMM | 1,358 TFLOPS | 343 GB/s | 1.2x |
| 4 groups x 8192, N 4096, K 7168 | Grouped (contiguous) | 1,297 TFLOPS | 418 GB/s | 1.2x |
| 1 group x 1024, N 4096, K 7168 | Grouped (masked) | 1,233 TFLOPS | 924 GB/s | 1.2x |
Selected rows from the launch README's H800 SXM5 benchmark table, as DeepSeek published them on 26 February 2025 [16].
These are DeepSeek's own numbers and depend heavily on matrix dimensions, GPU, and compiler version, so they are best read as representative rather than universal. They also moved over time. An updated table on 25 March 2025 raised the 4096x7168x16384 dense figure to 1,426 TFLOPS and the contiguous grouped figures to 1,346 TFLOPS at 1.3x [31]. A news entry dated 18 April 2025 reported DeepGEMM "now achieves up to 1550 TFLOPS on H800" [1]. The Performance section was removed from the README on 14 April 2025, in the commit that replaced output padding with shared-memory swizzling, and as of 1 October 2026 the repository publishes no benchmark table for the CUDA library [1][16].
| Property | Detail |
|---|---|
| Developer | DeepSeek |
| Released | 26 February 2025 (Open Source Week, day 3) |
| Repository created | 13 February 2025 |
| Target hardware | NVIDIA SM90 (Hopper, including H100 and H800) and SM100 (Blackwell) |
| Precision | FP8, FP4 and BF16, with fine-grained scaling and higher-precision second accumulation |
| Operations | Dense GEMM, grouped GEMM (contiguous, masked, K-grouped), fused Mega MoE, MQA indexer scoring, HyperConnection |
| Compilation | Just-in-time through DeepJIT, no CUDA build at installation |
| Requirements (Oct 2026) | CUDA Toolkit 12.9 or higher, Python 3.8 or higher, PyTorch 2.3 or higher, CUTLASS 4.0 or higher, C++20 |
| Latest public release | 30 September 2026 (PR #462) |
| Package version | 2.8.1 |
| GitHub stars | 7,906 (1 October 2026) |
| License | MIT |
| Ascend port | DeepGEMM-Ascend, released 30 September 2026 |
Sources: the repository README, the GitHub API and the installed package metadata [1][18][28].
How does DeepGEMM relate to CUTLASS and DeepSeek's models?
DeepGEMM describes itself as drawing on ideas from NVIDIA's CUTLASS and its CuTe layout library, while deliberately avoiding a heavy dependency on them. The repository puts it plainly: it "leverages some concepts from CUTLASS and CuTe, but avoids heavy reliance on their templates or algebras" [1]. CUTLASS is a large, general, template-heavy framework that covers many data types and architectures; DeepGEMM trades that generality for a small set of readable kernels [1][5]. The launch benchmark comparisons were made against an internal CUTLASS 3.6 implementation, which is the natural baseline for this kind of work [16].
Inside DeepSeek's own stack, DeepGEMM is the matrix-multiply engine behind both training and inference of the company's models, starting with V3 and R1 [1][2]. The same fine-grained quantization and promotion techniques that the V3 technical report describes are what the public kernels implement, and the 2025 hardware paper says DeepSeek's "fine-grained FP8 GEMM implementation has been open-sourced in DeepGEMM" [3][13]. When DeepSeek shipped V3.2-Exp in September 2025 with its sparse lightning-indexer attention, DeepGEMM gained scoring kernels for that indexer, described in the README as weighted-ReLU MQA logits, on 28 September 2025, and serving frameworks such as vLLM and SGLang wired those kernels in for day-zero support [1][7][11]. The indexer family has since split into non-paged kernels for prefilling and paged kernels for decoding, with MXFP4 and MXFP8 inputs carrying packed UE8M0 scales on SM100 [1].
DeepGEMM also sits alongside FlashMLA and DeepEP in a broader pattern: DeepSeek releases not only model weights but the low-level systems software needed to train and serve them. By 2026 those siblings had been joined by DeepSelect for top-k selection, TileKernels for vector and memory kernels written in TileLang, and DeepJIT for runtime compilation [19][22][23].
How has DeepGEMM evolved since launch?
DeepGEMM did not stop at the launch feature set. The table below follows the News section of the README, which is DeepSeek's own record of the public releases [1][16].
| Date | Change |
|---|---|
| 18 April 2025 | Reported up to 1,550 TFLOPS on H800 |
| 7 May 2025 | NVRTC compilation path, reported as up to 10x faster to compile (later disabled) |
| 14 May 2025 | Weight gradient kernels for dense and MoE backward passes |
| 20 July 2025 | Support for both SM90 and SM100, full refactor with a low-CPU-overhead JIT C++ module; FFMA SASS post-optimization retired |
| 28 September 2025 | Scoring kernels (weighted ReLU MQA logits) for the DeepSeek V3.2 lightning indexer |
| 16 April 2026 | Mega MoE, FP8xFP4 GEMM, FP4 indexer, Programmatic Dependent Launch (PDL), faster JIT compilation |
| 10 September 2026 | Sparse Indexer, Mega Gate, Mega mHC, DeepJIT, MoE and indexer optimizations |
| 30 September 2026 | DeepGEMM-Ascend released; further optimizations including locality-domain features |
The July 2025 refactor was the architectural turning point. It widened the supported memory layouts: where the SM90 kernels handle only the NT layout (row-major times column-major) that DeepSeek's models use, the SM100 kernels support all four combinations, NT, TN, NN and TT [1]. It also changed the scaling-factor format between architectures, with SM90 taking FP32 scales and SM100 taking packed UE8M0 [1].
The April 2026 release added Mega MoE, which the README describes as fusing and overlapping "EP dispatch, linear 1 and linear 2 (FP8xFP4 or FP8xFP8), SwiGLU, and EP combine into a single mega-kernel, overlapping NVLink communication and tensor core computation" [1]. It requires a multi-process launch with symmetric memory and PyTorch 2.9 or newer, and the API exposes helpers to allocate the symmetric buffer, transform expert weights into the required layout, and optionally localize them into locality domains [1]. The fused path covers the whole expert computation: dispatch, the first linear layer, the SwiGLU activation, the second linear layer, and the combine. The FP4 work connects DeepGEMM to the microscaling formats and NVFP4 directions that hardware vendors are pushing for lower-precision inference.
By version 2.8.1 the exported API spans considerably more than GEMM. It includes FP8 and FP4 GEMMs in all four layouts with M-grouped and K-grouped variants, BF16 GEMMs, a BF16 Mega Gate kernel, einsum kernels, dense and sparse MQA logit kernels in paged and non-paged forms, a mega_mhc kernel and a TensorFloat-32 HyperConnection pre-normalization GEMM, locality-domain helpers, thin wrappers over cuBLASLt, and a set of legacy Triton kernels for A100 that the module header marks as possibly deprecated "or rewrite in TileLang" [18]. Through these changes the design philosophy stayed constant: keep the core small, compile at runtime, and lean on a handful of well-understood hardware features rather than a large abstraction layer.
DeepGEMM-Ascend
DeepSeek published DeepGEMM-Ascend on 30 September 2026. The repository was created on 29 September 2026, carries a single public commit dated 30 September 2026, and is released under the MIT License with a 2026 DeepSeek copyright line; it had 421 stars and 29 forks on 1 October 2026, and the Python package reports version 0.1.0 [19][20][21].
The README calls it "a port of DeepGEMM to the HUAWEI Ascend platform" that "is fully API-compatible with DeepGEMM and supports BF16, FP8, FP4 GEMM, MQA logits, and MegaMoE." Because it keeps the same package name, deep_gemm, DeepSeek says users "can simply install the package and use the same APIs and development workflow as DeepGEMM on other supported platforms" [19]. The stated design idea is a thin layer over the Ascend matrix multiply-add (MAD) primitives that hides "fractal layouts, alignment constraints, address calculations, and verbose low-level parameters," and the project says it uses Ascend-specific techniques including "sparse data loading and coroutine-based pipelining, to approach the performance limits of the Ascend hardware" [19].
Hardware and software requirements
| Requirement | Detail |
|---|---|
| Hardware | Huawei Ascend NPU, "developed and validated on the Ascend 950 series" |
| Toolkit | CANN 9.20, supplying bin/bisheng and bin/ld.lld |
| Framework | torch_npu |
| Python | 3.10 or higher |
| Compiler | C++20 <format> support |
| Declared dependency | tilelang, used by the HC prenorm kernel |
| Build-from-source extras | tree-sitter and tree-sitter-cpp, for generating Python type stubs |
| Submodules | DeepJIT (JIT runtime), clangd-ascend, magic_enum |
| License | MIT |
Source: the DeepGEMM-Ascend README, .gitmodules and setup.py [19].
The port links against Ascend libraries including ascendcl, hccl_fwk, opapi and torch_npu, and finds the toolkit through ASCEND_HOME_PATH or ASCEND_TOOLKIT_HOME [19]. Two differences from the CUDA version are called out explicitly. The scaling-factor format diverges: 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 optimal hardware efficiency" [19]. And several utilities are Ascend-specific, including set_npu_arch and get_npu_arch to override or query the dav-* architecture used for JIT compilation, a set_num_sms that sets the number of AI cores rather than streaming multiprocessors, and a family of ACLNN reference GEMMs used to cross-check the kernels in tests [19].
The exported API is also narrower than upstream. The Ascend package exposes dense and contiguous M-grouped and K-grouped GEMMs in FP8, FP4 and BF16, einsum kernels, paged and non-paged MQA logits, layout transforms, locality-domain helpers and Mega MoE, but not the masked grouped GEMM, the sparse indexer kernels, Mega Gate or mega_mhc that the CUDA package carries [18][19]. The HyperConnection pre-normalization kernel is implemented in TileLang rather than in the port's own kernel code, and the acknowledgements credit TileLang for backing the mHC kernel and DeepJIT for the JIT runtime [19].
Reported performance
DeepSeek states that "despite its lightweight codebase, DeepGEMM Ascend can achieve peak hardware performance across a wide range of matrix shapes," and the release note says "the kernels are designed to achieve near-peak hardware performance" [19]. These are the project's own claims. The published figures were "measured on Ascend 950DT (CANN 9.20) using bench_msprof with cold L2," on shapes and cases that "follow the DeepGEMM test suite and cover inference and training workloads of the DeepSeek model series" [19]. Covering the release, the trade publication TNGlobal wrote that the announcement "did not provide independent performance results or disclose production deployments using the full stack, so comparisons with established accelerator platforms remain unverified" [24].
For dense GEMM at 4096x7168x16384, DeepSeek reports the following, saying the kernels reach "up to 99.8% of the hardware limit across types" [19]:
| Type | Latency | Computation | Stated hardware limit | Utilization |
|---|---|---|---|---|
| FP4xFP4 | 565.6 us | 1,701 TFLOPS | 1,730 TFLOPS | 98.3% |
| FP8xFP4 | 1,117.8 us | 861 TFLOPS | 865 TFLOPS | 99.5% |
| FP8xFP8 | 1,117.6 us | 861 TFLOPS | 865 TFLOPS | 99.5% |
| BF16xBF16 | 2,229.9 us | 431 TFLOPS | 432 TFLOPS | 99.8% |
Other reported results, all from the same README and the same Ascend 950DT setup [19]:
- M-grouped GEMM for DeepSeek MoE experts (FP8xFP4 with BF16 output) reaches 862 TFLOPS at eight groups of 4,096 rows with N 7168 and K 3072.
- MQA logits for the lightning indexer top out at 743 TFLOPS in prefill and 708 TFLOPS in decoding, both with FP4 inputs. DeepSeek notes the kernel is "FIX-pipe bound rather than compute bound, saturating the FIX pipe at 99% utilization."
- Mega MoE, benchmarked over EP8 with top-k 6 and one shared expert and averaged across eight ranks, reaches 846.3 TFLOPS at 384 experts, hidden size 7168, intermediate size 3072 and 16,384 tokens.
- The HC prenorm GEMM for what DeepSeek calls the mHC (Manifold-Constrained Hyper-Connections) module "nearly saturates the HBM write bandwidth," reaching 3,463 GB/s at M 8192, N 24 and K 28672.
The README names Kexing Zhou, Zhean Xu and Chenggang Zhao as project leads, lists per-component contributors, and thanks Huawei "for its technical support and engineering expertise throughout the development of DeepGEMM-Ascend" [19].
The 30 September 2026 Ascend release
DeepGEMM-Ascend arrived as one piece of a larger announcement. Reuters reported on 30 September 2026 that DeepSeek said in a post on its official WeChat account that it had partnered with Huawei Technologies to develop programming tools optimised for Ascend chips, and was open-sourcing programming infrastructure for the platform including compute and communication libraries [22]. Reuters framed the move as "the latest sign of Chinese tech firms deepening ties as they seek alternatives to the Nvidia ecosystem," and noted the announcement came two weeks after Huawei unveiled its next generation of AI processors and supernode computing systems [22]. According to Reuters, DeepSeek said Huawei "provided full support" in developing the infrastructure, and the two companies "jointly advanced a 'supernode' solution based on 128 Ascend 950 chips, optimising both computation and communication" [22].
The Chinese outlet IT之家 (IT Home), reporting the same WeChat post, listed six components and the role DeepSeek assigned to each: TileLang as the high-level language and compilation tool, DeepGEMM to accelerate general matrix operations, DeepEP for large-scale cross-device communication, TileKernels for routine vector-compute and memory-access operators, FlashMLA for sparse attention to improve long-context efficiency, and DeepSelect for efficient data filtering [23]. IT Home reported DeepSeek's statement that all the components correspond one to one with its earlier releases for the NVIDIA platform, that the Ascend version of TileLang wraps Ascend C low-level instructions in a high-level programming interface without losing hardware performance, and that in multiple key test cases the components' compute and communication performance was already close to the hardware limit [23]. DeepSeek also said the TileLang approach was validated first on NVIDIA hardware and now carries most of the operator implementations in DeepSeek-V4 series training [23].
On TileLang itself, Reuters quoted DeepSeek: "To build a new generation of independent, self-controlled GPU software ecosystems, the first priority is establishing a high-level language that is universal, easy to program, and still capable of reaching the hardware's full performance potential," and "TileLang was created precisely to meet this need," offering "a simpler programming model" than NVIDIA's CUDA [22]. TileLang is not a DeepSeek project: it is the tile-ai project, originally developed by researchers at Peking University and Microsoft Research, and DeepSeek has been using it for about a year, according to The Decoder's account of the release [25].
The sibling repositories carry matching release notes dated 30 September 2026. DeepEP-Ascend, created on 30 September 2026, provides expert-parallel all-to-all dispatch and combine for Ascend NPUs with buffer APIs aligned to the NVIDIA version, with kernels written in Ascend C over HCCL/HCOMM, UBMEM and URMA and compiled at runtime through DeepJIT; its own measurements were taken on Ascend 950DT with CANN 9.2.0 [26]. FlashMLA released sparse attention prefill and decoding kernels for the Ascend 950 NPU, reporting up to 410 TFLOPS in prefill and 360 TFLOPS in decoding, which it puts at 95 percent and 83 percent of hardware peak, and in the same release removed support for the NVIDIA Hopper architecture and for DeepSeek V3, V3.2 and V4.0 [27]. 智东西 (Zhidongxi), a Chinese technology outlet, summarised the release on X the same day as DeepSeek open-sourcing its AI infrastructure stack for Ascend NPUs, listing the same six components [30].
Who uses DeepGEMM, and what are its limitations?
After release DeepGEMM was studied closely by the open-source community as a compact, documented example of production FP8 kernels, and the repository accumulated thousands of GitHub stars within its first year [1][28]. Inference frameworks adopted it as a building block, most visibly in the day-zero support that vLLM and SGLang built for DeepSeek-V3.2's sparse attention using DeepGEMM's indexer kernels [1][11].
The main limitation has always been hardware scope. The CUDA library requires an NVIDIA SM90 or SM100 GPU and uses architecture-specific features such as the TMA, so its kernels do not run on other vendors' accelerators; the repository does ship a small set of legacy Triton kernels for A100, but the module header marks them as candidates for deprecation [1][18]. Rather than generalise the CUDA codebase, DeepSeek answered the portability question with a separate repository, DeepGEMM-Ascend, which reimplements the same APIs against Ascend primitives [19]. The published benchmark numbers for both versions are DeepSeek's own and tied to specific shapes, chips and toolkit versions, so they should be treated as indicative rather than guaranteed for an arbitrary workload, and the CUDA README no longer publishes a benchmark table at all [1][16][19]. Because the library specialises aggressively to the shapes it sees, the first call for a new shape also pays a JIT compilation cost before the fast kernel is available, although caching through DeepJIT reduces that overhead in practice [1][29].
References
- ^1 ^2 ^3 ^4 ^5 ^6 ^7 ^8 ^9 ^10 ^11 ^12 ^13 ^14 ^15 ^16 ^17 ^18 ^19 ^20 ^21 ^22 ^23 ^24 ^25 ^26 ^27 ^28 ^29 ^30 ^31 ^32 ^33 ^34 ^35 ^36 ^37DeepSeek. "DeepGEMM: clean and efficient BLAS kernel library on GPU." GitHub repository README, 2025-2026. github.com/...DeepGEMM
- ^1 ^2 ^3 ^4 ^5DeepSeek. "open-infra-index: Open Source Week." GitHub, 2025. github.com/...open-infra-index
- ^1 ^2 ^3 ^4 ^5 ^6 ^7 ^8 ^9 ^10 ^11DeepSeek-AI. "DeepSeek-V3 Technical Report." arXiv preprint arXiv:2412.19437, 2024. arxiv.org/...2412.19437
- ^MarkTechPost. "DeepSeek AI Releases DeepGEMM: An FP8 GEMM Library that Supports both Dense and MoE GEMMs Powering V3/R1 Training and Inference." 25 February 2025. marktechpost.com/...g-v3-r1-training-and-inference
- ^NVIDIA. "CUTLASS: CUDA Templates for Linear Algebra Subroutines." GitHub repository. github.com/...cutlass
- ^NVIDIA. "NVIDIA Hopper Architecture." NVIDIA Data Center. nvidia.com/...hopper-architecture
- ^DeepSeek-AI. "DeepSeek-V3.2-Exp." GitHub repository, 2025. github.com/...DeepSeek-V3.2-Exp
- ^1 ^2Micikevicius, Paulius, et al. "FP8 Formats for Deep Learning." arXiv preprint arXiv:2209.05433, 2022. arxiv.org/...2209.05433
- ^DigiAlps. "DeepSeek AI Drops DeepGEMM, An FP8 GEMM Library That Powers V3 and R1 AI Models." 2025. digialps.com/...ry-that-powers-v3-and-r1-ai-models
- ^DeepSeek. "DeepEP: an efficient expert-parallel communication library." GitHub, 2025. github.com/...DeepEP
- ^1 ^2vLLM. "DeepSeek-V3.2-Exp in vLLM: Fine-Grained Sparse Attention in Action." vLLM Blog, 29 September 2025. blog.vllm.ai/...deepseek-v3-2
- ^Zhang, Jinpeng. "DeepSeek Technical Analysis (5): FP8 Training." Medium, 2025. dataturbo.medium.com/...-fp8-training-ff34768727b8
- ^1 ^2DeepSeek-AI. "Insights into DeepSeek-V3: Scaling Challenges and Reflections on Hardware for AI Architectures." arXiv preprint arXiv:2505.09343, 2025. arxiv.org/...2505.09343
- ^NVIDIA. "NVIDIA Hopper Tuning Guide" (Tensor Memory Accelerator section). NVIDIA CUDA documentation. docs.nvidia.com/...hopper-tuning-guide
- ^DeepSeek. "Day 3 of #OpenSourceWeek: DeepGEMM." X (formerly Twitter), 26 February 2025. x.com/...1894553164235640933
- ^1 ^2 ^3 ^4 ^5 ^6 ^7 ^8 ^9 ^10 ^11 ^12 ^13 ^14 ^15DeepSeek. DeepGEMM README at the launch commit (78cacf70d4), 26 February 2025. GitHub. github.com/...README.md
- ^DeepSeek. "DeepGEMM Scaling Factor Format," docs/scaling-factor-format.md. GitHub. github.com/...scaling-factor-format.md
- ^1 ^2 ^3 ^4 ^5DeepSeek. `deep_gemm/__init__.py` (version 2.8.1), deepseek-ai/DeepGEMM. GitHub. github.com/...__init__.py
- ^1 ^2 ^3 ^4 ^5 ^6 ^7 ^8 ^9 ^10 ^11 ^12 ^13 ^14 ^15 ^16 ^17 ^18DeepSeek. "DeepGEMM Ascend" README, plus `.gitmodules`, `setup.py` and `deep_gemm/__init__.py`. GitHub, deepseek-ai/DeepGEMM-Ascend, 30 September 2026. github.com/...DeepGEMM-Ascend
- ^GitHub REST API. "deepseek-ai/DeepGEMM-Ascend" repository metadata (created_at, description, license, stars, forks). Accessed 1 October 2026. api.github.com/...DeepGEMM-Ascend
- ^DeepSeek. LICENSE (MIT License), deepseek-ai/DeepGEMM-Ascend. GitHub. github.com/...LICENSE
- ^1 ^2 ^3 ^4 ^5 ^6Reuters. "DeepSeek partners with Huawei to develop chip programming tools." 30 September 2026 (via Business Standard). business-standard.com/...ming-tools-126093001470_1
- ^1 ^2 ^3 ^4Qingyuan. "DeepSeek open-sources infrastructure components for Huawei's Ascend computing platform, matching its NVIDIA components one to one." IT之家 (IT Home), 30 September 2026. ithome.com/...604
- ^TNGlobal. "DeepSeek and Huawei open Ascend AI programming tools." TechNode Global, 1 October 2026. technode.global/...wei-ascend-ai-programming-tools
- ^Schreiner, Maximilian. "China's AI industry closes ranks as Deepseek ships open-source software for Huawei's Ascend chips." The Decoder, 30 September 2026. the-decoder.com/...ftware-for-huaweis-ascend-chips
- ^DeepSeek. "DeepEP-Ascend" README. GitHub, deepseek-ai/DeepEP-Ascend. Accessed 1 October 2026. github.com/...DeepEP-Ascend
- ^DeepSeek. "FlashMLA" README (News, 2026.09.30 release). GitHub, deepseek-ai/FlashMLA. Accessed 1 October 2026. github.com/...FlashMLA
- ^1 ^2 ^3 ^4 ^5GitHub REST API. "deepseek-ai/DeepGEMM" repository metadata and README commit history. Accessed 1 October 2026. api.github.com/...DeepGEMM
- ^1 ^2 ^3GitHub REST API. "deepseek-ai/DeepJIT" repository metadata. Accessed 1 October 2026. api.github.com/...DeepJIT
- ^智东西 (Zhidongxi, @Chinazhidx). Post on DeepSeek's open-source Ascend stack. X, 30 September 2026. x.com/...2105138013609369641
- ^DeepSeek. DeepGEMM README at commit 55ab91f72f (updated performance table), 25 March 2025. GitHub. github.com/...README.md
- ^DeepSeek. DeepGEMM README at commit d75b218b7b (NVRTC news entry), 7 May 2025. GitHub. github.com/...README.md
Improve this article
Add missing citations, update stale details, or suggest a clearer explanation. Every suggestion is reviewed for sourcing before it goes live.
8 revisions · v8 · 5,305 words · full history
Fact-checks are independent of edits: a reviewer re-verifies the article against its sources and stamps the date. How we verify
Research and drafting on this wiki are AI-assisted, under named human editorial standards. How AI is used here
Reviewer note: Independently fact-checked 1 Oct 2026 against the repo at named commits, the DeepSeek-V3 paper and DeepGEMM-Ascend; 5 defects corrected incl. when the README performance section was removed
Cite this page: AI Wiki. "DeepGEMM." aiwiki.ai, updated 1 Oct 2026, fact-checked 1 Oct 2026. CC BY 4.0. https://aiwiki.ai/wiki/deepgemm