# FlashMLA

> Source: https://aiwiki.ai/wiki/flashmla
> Updated: 2026-10-01
> Fact-checked: 2026-10-01
> Categories: AI Inference, AI Infrastructure, Developer Tools, Open Source AI
> License: CC BY 4.0 (https://creativecommons.org/licenses/by/4.0/) - attribute to "AI Wiki (aiwiki.ai)"
> Cite as: AI Wiki. "FlashMLA." aiwiki.ai, 1 Oct 2026. https://aiwiki.ai/wiki/flashmla
> From AI Wiki (https://aiwiki.ai), the free encyclopedia of artificial intelligence. Reuse freely with attribution.

**FlashMLA** is an open-source library of attention kernels from [DeepSeek](https://aiwiki.ai/wiki/deepseek), built around Multi-head Latent Attention (MLA), the attention variant DeepSeek uses to shrink the [KV cache](https://aiwiki.ai/wiki/kv_cache) in its large language models. DeepSeek published it on 24 February 2025 as the first project of its "Open Source Week," a run of five daily code drops pulled straight from the company's production inference stack. At that point it was one thing: a BF16 decoding kernel for NVIDIA [Hopper](https://aiwiki.ai/wiki/nvidia_hopper) GPUs with a paged KV cache, reporting up to 3000 GB/s in memory-bound settings and 580 TFLOPS in compute-bound settings on an H800. [1][2]

It has since become the kernel library behind [DeepSeek Sparse Attention](https://aiwiki.ai/wiki/deepseek_sparse_attention) (DSA). As of 1 October 2026 the README describes FlashMLA as "DeepSeek's library of optimized attention kernels, powering inference of the DeepSeek-V4.1 model on NVIDIA GPUs and Huawei Ascend NPUs," referring to [DeepSeek-V4.1-Flash](https://aiwiki.ai/wiki/deepseek_v4_1_flash). The current code targets NVIDIA SM100 and SM103 ([Blackwell](https://aiwiki.ai/wiki/nvidia_blackwell)-generation) GPUs and the [Huawei Ascend](https://aiwiki.ai/wiki/huawei_ascend) 950 NPU; a release on 30 September 2026 removed support for Hopper and for the earlier DeepSeek V3, V3.2 and V4.0 models, and changed the FP8 and FP4 KV cache format, so that version is not compatible with previous ones. The repository is MIT-licensed and had about 13,000 GitHub stars on 1 October 2026. [1]

The project sits at the intersection of two trends in efficient inference. One is the attention redesign DeepSeek introduced with [DeepSeek-V2](https://aiwiki.ai/wiki/deepseek_v2) to cut the memory cost of long contexts. The other is the family of fused attention kernels, led by [FlashAttention](https://aiwiki.ai/wiki/flash_attention), that keep the GPU busy by avoiding round trips to slow high-bandwidth memory. FlashMLA applies the second idea to the first, for the specific tensor shapes MLA produces. [1][4]

## What is FlashMLA?

FlashMLA is not a model and not a serving framework; it is the set of low-level routines a serving stack calls to compute attention. The current repository contains four groups of kernels, all described in the README: [1]

| Kernel group | What it does |
|---|---|
| Sparse prefill | Token-level sparse attention for the prefill stage |
| Sparse decoding | Token-level sparse attention for decoding, reading an FP8 or FP4 KV cache |
| Fused norm + RoPE + attn + RoPE + cast | One kernel combining Q-norm, Q-RoPE, core attention, O-RoPE (conjugate) and the cast to FP8 |
| Dense attention | Dense multi-head attention for prefill, plus the backward pass |

The sparse kernels are the ones that power DSA, the mechanism DeepSeek introduced with DeepSeek-V3.2-Exp. The README points to the DeepSeek-V3.2-Exp repository for the DSA description. [1]

The day-one release was narrower. DeepSeek's Open Source Week index billed it as an "Efficient MLA Decoding Kernel for Hopper GPUs," "Optimized for variable-length sequences, battle-tested in production," with BF16 support, a paged KV cache with block size 64, and "3000 GB/s memory-bound | BF16 580 TFLOPS compute-bound on H800." The phrase about production matters: FlashMLA was not a research prototype but code taken from the stack that served DeepSeek's own chat and API traffic, which means it had to cope with requests arriving at different lengths and batch sizes shifting from moment to moment. [2]

The README credits the kernel's lineage directly: FlashMLA "is inspired by FlashAttention 2&3 and cutlass projects." In short, it borrows FlashAttention's fused, tiled approach to attention and [CUTLASS](https://aiwiki.ai/wiki/cutlass)'s tensor-core scheduling, then specializes both for MLA. [1]

## When was FlashMLA released, and what was Open Source Week?

DeepSeek created the FlashMLA repository on 21 February 2025 and announced Open Source Week the same day, saying the drops would begin on 24 February. The company framed the effort plainly. Its index repository describes the team as "a tiny team @deepseek-ai pushing our limits in AGI exploration" and says the releases are "humble building blocks of our online service: documented, deployed, and battle-tested in production. No vaporware, just sincere code that moved our tiny yet ambitious dream forward." Over the week DeepSeek shipped FlashMLA on day one, the [DeepEP](https://aiwiki.ai/wiki/deepep) expert-parallel communication library on day two, the [DeepGEMM](https://aiwiki.ai/wiki/deepgemm) FP8 matrix-multiply library on day three, a set of parallelism tools including [DualPipe](https://aiwiki.ai/wiki/dualpipe) and EPLB on day four, and the Fire-Flyer File System on day five. A sixth post, billed as "One More Thing," described the DeepSeek-V3/R1 inference system. [1][2]

| Day | Date (2025) | Project | What it does |
|---|---|---|---|
| 1 | 24 Feb | FlashMLA | MLA decoding kernel for Hopper GPUs |
| 2 | 25 Feb | DeepEP | Expert-parallel communication library for MoE |
| 3 | 26 Feb | DeepGEMM | FP8 general matrix-multiply library |
| 4 | 27 Feb | DualPipe / EPLB | Pipeline parallelism and expert load balancing |
| 5 | 28 Feb | 3FS (Fire-Flyer File System) | High-throughput distributed file system |

## What is MLA, and how does FlashMLA relate to it?

MLA itself predates the kernel. DeepSeek introduced [Multi-head Latent Attention](https://aiwiki.ai/wiki/multi-head_latent_attention) in the DeepSeek-V2 technical report in May 2024 and carried it forward into [DeepSeek-V3](https://aiwiki.ai/wiki/deepseek_v3) and the R1 reasoning model. The motivation is the KV cache. During autoregressive generation a transformer stores the key and value vectors for every past token so it does not recompute them at each step. That cache grows with sequence length and [batch size](https://aiwiki.ai/wiki/batch_size), and on long contexts it can dominate GPU memory, which caps how many requests a server can run at once. [4][6]

The distinction is worth stating plainly: MLA is the attention architecture baked into the model weights, while FlashMLA is the kernel code that executes that architecture efficiently. You can define a model with MLA and run it with a generic attention kernel; FlashMLA is the purpose-built routine that makes MLA fast on the hardware DeepSeek serves on. [1][4]

## How does MLA differ from MHA and GQA?

Standard multi-head attention gives every attention head its own keys and values, so the cache holds a full set of key and value vectors per head per token. [Grouped-query attention](https://aiwiki.ai/wiki/grouped_query_attention) reduces that by letting several query heads share one key-value head, which is the path used by models such as [Llama 2](https://aiwiki.ai/wiki/llama_2) 70B and many others. Multi-query attention, proposed by Noam Shazeer in 2019, pushes the same idea to a single shared key-value head. Both trade a little quality for a smaller cache. [4][9]

MLA takes a different route. Instead of cutting the number of key-value heads, it compresses the keys and values themselves. The DeepSeek-V2 paper describes a low-rank joint compression that projects the hidden state down into a single small latent vector. Only that latent vector goes into the cache. When the model needs the actual keys and values for an attention step, it reconstructs them from the latent vector through up-projection matrices. The cache therefore stores one compact representation per token rather than a full set of per-head key and value vectors. [4]

Rotary position embedding does not fit this scheme cleanly, because the position-dependent rotation has to be applied before the projections in a way that breaks the compression trick. DeepSeek handled this with a decoupled RoPE design that carries the rotary signal on extra dedicated query and key dimensions, kept separate from the compressed part. The result keeps positional information without forcing the model to cache full keys and values. [4]

The payoff is large. The DeepSeek-V2 paper reports that, compared with DeepSeek 67B, V2 "saves 42.5% of training costs, reduces the KV cache by 93.3%, and boosts the maximum generation throughput to 5.76 times." The cache figure is the MLA result; the throughput figure is for the whole model, which also replaced dense feed-forward layers with DeepSeekMoE. Both describe the architecture, not the FlashMLA kernel. [4]

| Attention variant | What gets cached | Relative KV cache | Used by |
|---|---|---|---|
| Multi-head attention (MHA) | Full keys and values per head | Largest | GPT-style models, DeepSeek 67B |
| Grouped-query attention (GQA) | Keys and values per head group | Reduced | Llama 2 70B and many others |
| Multi-query attention (MQA) | One shared key-value head | Small | Introduced by Shazeer (2019) [9] |
| Multi-head latent attention (MLA) | One compressed latent vector | Smallest reported | DeepSeek-V2, DeepSeek-V3, DeepSeek-R1 |

The table is a qualitative summary. Exact ratios depend on head counts, dimensions, and the compression rank, and they vary by model.

## Why does MLA need its own kernel?

The compression that makes MLA cheap on memory also changes the math the GPU has to run. The cached latent vector has to be expanded back into keys and values, and the attention computation works over tensor shapes that differ from plain multi-head or grouped-query attention. A generic attention kernel can run MLA, but it will not lay out the work in the way that keeps tensor cores and the memory pipeline near their limits. A purpose-built kernel can.

A common assumption is that decode-stage attention is always memory-bound, because the model generates one token per step per sequence and so has very little compute per byte touched. DeepSeek's own analysis says MLA decode is the exception. The April 2025 deep-dive derives the kernel's compute-to-memory ratio as roughly two times the number of query heads times the number of query tokens per request, and notes that an H800 SXM5 has a peak memory bandwidth of 3.35 TB/s and peak throughput of 990 TFLOPS, dropping to about 865 TFLOPS in practice because of clock throttling to around 1600 MHz. On those numbers the kernel is compute-bound once query heads times query tokens reaches 128. Because DeepSeek does not use tensor parallelism for its decoding instances, the query-head count is 128, which the document says puts the kernel in the compute-bound regime. The optimization target for the 2025 rewrite was therefore compute throughput, not bandwidth. [3]

## How is FlashMLA designed?

The FlashAttention and CUTLASS influences map onto two parts of the design. From FlashAttention it borrows the fused, tiled approach to attention. Rather than writing the large intermediate attention-score matrix out to high-bandwidth memory and reading it back, the kernel streams blocks of keys and values through fast on-chip memory and computes the softmax in an online, running fashion. FlashAttention-3 in particular was written to use Hopper features, which the early FlashMLA also targeted. [1][7]

From CUTLASS it inherits the building blocks for high-performance matrix multiply on NVIDIA tensor cores. CUTLASS is NVIDIA's template library of [CUDA](https://aiwiki.ai/wiki/cuda) primitives for linear algebra, and it gives kernel authors a way to schedule tensor-core work, manage shared memory, and pipeline data movement at a level close to the metal. MLA attention is, underneath, a sequence of matrix multiplies tied together by the attention pattern, so CUTLASS-style scheduling is a natural fit. The dense prefill and backward kernels in the current repository are built on CUTLASS fused multi-head attention code, and the build script exempts them from the register-spill check it applies to FlashMLA's own kernels. [1][8]

Two serving features shaped the original design and survive in the current one. The paged KV cache stores the cache in fixed-size blocks rather than one contiguous buffer per sequence, the same idea popularized by PagedAttention in [vLLM](https://aiwiki.ai/wiki/vllm). Paging reduces memory fragmentation and lets a server pack more sequences into the same memory, which raises the achievable batch size. The variable-length support means the kernel handles a batch where each sequence has a different length without padding everything to the longest one. [1]

### Requirements and build

The README lists separate requirements per platform. As of 1 October 2026: [1]

| Platform | Hardware | Toolchain |
|---|---|---|
| NVIDIA | SM100 or SM103 GPU | CUDA 13.1 and above (13.2 and above recommended, for the native `cvt.rn.bf16x2` conversions the quantized decoding kernels use), PyTorch 2.0 and above |
| Huawei | Ascend 950 NPU | CANN 9.2.0 and above (which supplies the `bisheng` compiler and the Ascend C headers), `torch_npu`, PyTorch 2.0 and above |

The build target is detected automatically: the presence of `/dev/davinci_manager` means Ascend, anything else means CUDA, and the choice can be forced with the `FLASH_MLA_BUILD_TARGET_PLATFORM` environment variable. The CUDA path compiles for the architecture-specific targets `sm_100a` and `sm_103a` rather than the family target `sm_100f`, because the kernels use a scaled `cvt` instruction; the Ascend path defaults to the NPU architecture `dav-3510` and compiles `.asc` Ascend C sources with `bisheng`. After a CUDA build the setup script scans the produced shared object and fails the build if any kernel the repository defines spills registers. [1][16]

## How fast is FlashMLA?

Every figure below is DeepSeek's own measurement as stated in the README or its deep-dive documents, on DeepSeek's own hardware, for a favorable configuration. Real throughput depends on model dimensions, batch size, sequence length and the accelerator in use.

| Kernel | Hardware and toolchain | Reported figure | First stated |
|---|---|---|---|
| Dense MLA decoding | H800 SXM5, CUDA 12.6 | up to 3000 GB/s memory-bound, 580 TFLOPS compute-bound | Feb 2025 |
| Dense MLA decoding | H800 SXM5, CUDA 12.8 | up to 660 TFLOPS compute-bound (a 5 to 15 percent gain) | Apr 2025 |
| Sparse prefill and decoding (DSA) | Hardware not stated in the README; the accompanying deep-dive is on H800 | up to 640 TFLOPS prefill, 410 TFLOPS decoding | Sep 2025 |
| Fused norm + RoPE + attn + RoPE + cast | B200, CUDA 13.3 | up to 1460 TFLOPS prefill, 950 TFLOPS decoding | Sep 2026 |
| Sparse prefill | B200, CUDA 13.3 | up to 1350 TFLOPS | Sep 2026 |
| Sparse decoding | B200, CUDA 13.3 | up to 1024 TFLOPS | Sep 2026 |
| Sparse prefill | Huawei Ascend 950 NPU | up to 410 TFLOPS, which DeepSeek puts at 95 percent of theoretical hardware peak | Sep 2026 |
| Sparse decoding | Huawei Ascend 950 NPU | up to 360 TFLOPS, or 83 percent of theoretical peak | Sep 2026 |
| Dense MHA prefill | B200 | up to 1460 TFLOPS forward and 1000 TFLOPS backward, "as reported by NVIDIA" | 2025 |

The original 580 TFLOPS and 3000 GB/s pair are the day-one numbers and were measured with CUDA 12.6; DeepSeek raised the recommended CUDA version to 12.8 the next day and restated the same throughput there. For the current release the README recommends the fused norm-RoPE-attn-RoPE-cast kernel over the standalone sparse prefill and decode kernels on CUDA, even though the fused kernel does more work, because it removes the launch overhead of the small operations around attention. [1][3]

## What changed after the initial release?

| Date | Change |
|---|---|
| 24 Feb 2025 | Day-one release: BF16 MLA decoding kernel for Hopper, paged KV cache with block size 64, CUDA 12.3 and above [1][2] |
| 24 to 25 Feb 2025 | FP16 support merged; recommended CUDA raised to 12.8 [1] |
| 22 Apr 2025 | Performance update giving 5 to 15 percent on compute-bound workloads, up to 660 TFLOPS on H800 SXM5, with an interface compatible with the old one, plus a deep-dive write-up [1][3] |
| 1 Aug 2025 | MHA forward and backward kernels for SM100, contributed by NVIDIA in pull request 76 [1] |
| 14 Aug 2025 | Support for CUDA below 12.8 dropped [1] |
| 29 Sep 2025 | Token-level sparse attention kernels for DSA, released with DeepSeek-V3.2, with an FP8 KV cache and a deep-dive on the Hopper FP8 sparse decoding kernel [1][5] |
| 10 Sep 2026 | Kernels for DeepSeek-V4.1: prefill and decoding with an FP8 or FP4 KV cache, plus the fused norm-RoPE-attn-RoPE-cast kernel [1][17] |
| 30 Sep 2026 | Sparse prefill and decoding kernels for the Huawei Ascend 950 NPU, a 10 to 15 percent decoding speed-up for the fused kernel, removal of Hopper and of V3/V3.2/V4.0 support, and a changed FP8/FP4 KV cache format [1][12][13] |

Users who need the older models or the old cache format are directed to a pinned commit from 15 September 2026 rather than to a release tag; the repository publishes no tags or releases. [1]

### The FP8 and FP4 KV cache

The decoding kernels now accept only quantized caches: unquantized bfloat16 is not supported. For DeepSeek-V4.1, whose head dimension is 512, the format is detected from the bytes per token in the cache tensor. In both layouts each token stores its quantized data first and its scales immediately after, a change from the earlier arrangement that grouped all of a block's data before all of its scales. [1]

| Layout | Bytes per token | Contents |
|---|---|---|
| V4.1 | 528 | 512 `float8_e4m3` values covering all 512 dimensions, including the 64 RoPE dimensions, then 16 `float8_e8m0` scales, one per 32 consecutive values |
| V4.1 FP4 | 288 | 512 `e2m1` values packed two per byte into 256 bytes, then 32 `float8_e4m3` scales, one per 16 consecutive values |

The earlier Hopper format introduced with DeepSeek-V3.2 was different in kind: tile-level quantization over the first 512 elements produced 512 `float8_e4m3` values and 4 `float32` scales, while the 64 RoPE elements were left in bfloat16 because they are sensitive to precision loss, for 656 bytes per token. DeepSeek's reason for adopting an FP8 cache at all was that V3.2 doubled the context window from 64K to 128K tokens, which it calculated would otherwise need 8.72 GiB of KV cache for a single 128K-token request. [5]

### Interfaces

Sparse decoding is driven by `flash_mla_with_kvcache`, preceded by a `get_mla_metadata` call that in the current release is a placeholder: the scheduling metadata is now generated on the first call to the kernel. Sparsity is expressed through an `indices` tensor of shape `(batch_size, seq_len_q, topk)` whose entries already encode the page block and the offset within it, so the kernel ignores the `block_table` argument it still requires positionally; invalid entries are set to -1. Sparse prefill uses `flash_mla_sparse_fwd` and takes optional attention-sink and per-token top-k-length arguments, but has no batch dimension, so multi-batch inference has to be simulated by reshaping. The dense kernels are exposed through `flash_attn_varlen_func` and two packed variants, which differ from the `flash_attn` package in returning both the output and the log-sum-exp. [1]

The fused kernel is the awkward one to adopt. It requires the Q_b and Wv projection weights to be permuted in advance, which the README shows doing with [TileKernels](https://aiwiki.ai/wiki/tilekernels)' `per_token_cast` quantizer and DeepGEMM's scale-factor layout transform, and its FP8 output is consumed directly by a DeepGEMM `fp8_einsum` against the permuted Wv weight. Running the fused kernel's test script requires three other DeepSeek-adjacent projects: [TileLang](https://aiwiki.ai/wiki/tilelang), TileKernels and DeepGEMM. [1]

## Does FlashMLA support Huawei Ascend?

Yes, since 30 September 2026, and in the same repository rather than a separate fork. On that day DeepSeek announced through its official WeChat account that it was open-sourcing infrastructure components for Huawei's Ascend computing platform: the TileLang high-level kernel language and its compiler support, plus compute and communication libraries, each corresponding to a component DeepSeek had already published for NVIDIA hardware. Chinese coverage of the announcement lists the six pieces as DeepGEMM for general matrix operations, DeepEP for large-scale cross-device communication, TileKernels for routine vector-compute and memory-access operators, FlashMLA for sparse attention operators that improve long-context processing, and [DeepSelect](https://aiwiki.ai/wiki/deepselect) for data filtering, alongside TileLang itself; the FlashMLA link given is the ordinary `deepseek-ai/FlashMLA` repository. DeepSeek said the components' compute and communication performance came close to hardware limits on several key test cases. [14][15]

In FlashMLA the Ascend work is a pull request, "Open-source release for Ascend 950 (2026.09.30)," merged on 30 September 2026, that touched 199 files. It adds sparse attention prefill and decoding kernels for the Ascend 950 NPU, reported at up to 410 TFLOPS in prefill (95 percent of the hardware's theoretical peak) and 360 TFLOPS in decoding (83 percent of peak), together with a technical report on the algorithms behind them. The kernels live under `csrc/ascend_kernels` as `.asc` Ascend C sources with six explicit instantiations covering prefill and decoding, with and without attention sinks, plus FP4 decoding variants. They are written against Ascend C directly and compiled with Huawei's `bisheng` compiler from CANN, not generated through TileLang, even though TileLang's Ascend port was the headline piece of the same announcement. [1][12][13][16]

Ascend coverage is partial. The README states that both the fused norm-RoPE-attn-RoPE-cast kernel and the dense MHA prefill kernel "currently support CUDA only; Ascend is not supported." Only the sparse prefill and decoding path exists for the NPU. [1]

### The Ascend kernel deep-dive

The accompanying technical report is unusually specific about why a kernel written for NVIDIA hardware does not transfer. Each AI Core of the Ascend 950 pairs one CUBE Core with two VECTOR Cores, and the authors explain that they could not keep to one VECTOR Core per CUBE Core for three reasons: a VECTOR Core can hold at most 16 issued-but-incomplete MTE2 copy requests at a time; dequantizing an FP8 or FP4 KV cache to bfloat16 cannot keep up with the CUBE Core at that ratio; and the FixPipe unit moves data out at only 128 bytes per cycle per AI Core, which would force a larger top-k tile size than the on-chip buffers can hold. The output accumulator for the online softmax therefore lives on the VECTOR Core, because data in the L0C buffer cannot be multiplied by a scalar in place. [12]

Several of the optimizations have no CUDA analogue. Because KV tokens are gathered sparsely, the usual on-the-fly ND-to-NZ layout conversion during the copy into L1 cannot be enabled, and the report notes that the AI Core's bus interface unit coalesces accesses only when source and destination are both contiguous, so the kernel routes KV through the unified buffer and converts the layout with a SIMD vector function instead. A "gather2" trick sets an MTE2 request's burst count to two so that two KV tokens move per request, easing the 16-request limit. The team also had to watch total code size, because the Ascend instruction caches are 32 KB on the CUBE Core, 16 KB on the VECTOR Core and 8 KB on the SIMD vector function. A "skip scale" optimization avoids rescaling the output accumulator while the running maximum stays within a threshold of the maximum currently used by the softmax, which on this hardware saves FixPipe bandwidth as well as arithmetic. The report also says the decoding KV cache layout was changed, interleaving each token's data with its scales, precisely because the previous arrangement forced two MTE2 copies per token on Ascend. [12]

### Reception

The South China Morning Post reported the release as "a step forward in the country's push to end dependence on processors from US giant Nvidia amid US export curbs," noting that DeepSeek released six software modules mirroring its earlier NVIDIA tools and aiming at an "independent and controllable" software ecosystem. The paper wrote that the Ascend work was made possible through "close collaboration" with Huawei, and that the two companies worked together on optimizing a supernode architecture powered by 128 Ascend 950 processors. [14]

## Ports by other vendors

The README's community support section points at third-party ports of FlashMLA to accelerators other than NVIDIA's and Huawei's. These are maintained outside DeepSeek's repository. [1]

| Vendor | Hardware | Port |
|---|---|---|
| MetaX | MetaX GPUs | MetaX-MACA/FlashMLA |
| Moore Threads | Moore Threads GPUs | MooreThreads/MT-flashMLA |
| Hygon | Hygon DCU | OpenDAS/MLAttention |
| Intellifusion | Intellifusion NNP | Intellifusion/tyllm |
| Iluvatar | Iluvatar Corex GPUs | Deep-Spark/FlashMLA, iluvatar branch |
| AMD | AMD Instinct | AITER/MLA in ROCm/aiter |

## How did the community adopt FlashMLA?

MLA support in open-source serving engines actually arrived before FlashMLA did. vLLM shipped MLA together with CUTLASS block FP8 kernels in version 0.7.1, announced in early February 2025, and reported a 3x increase in generation throughput and a 10x increase in token memory capacity against version 0.7.0, with token capacity on an 8x H200 machine rising from 54,560 to 512,000 and batch size from 13 to 128. The same write-up credits the "matrix absorption" algorithm from the DeepSeek-V2 paper, which computes directly on the latent cache rather than up-projecting it, and notes that multi-head attention still beat MLA on time-to-first-token at low query rates. [10]

[SGLang](https://aiwiki.ai/wiki/sglang) had gone first. Its v0.3 release in September 2024 claimed up to 7x higher throughput for DeepSeek MLA from weight absorption, grouped decoding kernels, FP8 batched matrix multiply and FP8 KV cache quantization, with benchmark results showing 3x to 7x over the baseline. [11]

Where FlashMLA mattered most was the sparse path. When DeepSeek-V3.2-Exp shipped in September 2025, vLLM's write-up said the engine "integrates with the lightning indexer CUDA kernels in DeepGEMM, as well as the new sparse attention kernel in FlashMLA," and that one reason it supports only a block size of 64 for the model is that "FlashMLA is tailored to it as well." Publishing the production kernel under a permissive license meant serving stacks could read the real code rather than reverse-engineer it from a paper. [1][18]

## Why does FlashMLA matter?

Inference cost for large models is dominated by how well a server keeps its accelerators full while holding the KV cache in memory, and MLA attacks the memory side hard. A kernel that runs MLA near hardware limits turns the architectural saving into a real serving saving, which is the difference between a clever idea on paper and lower cost per token in production. By open-sourcing the exact kernels it runs, DeepSeek let others reach a similar efficiency, which is part of why its models became practical to host outside DeepSeek itself. [1][10]

The release also fit a pattern in AI infrastructure where the hard-won engineering, not just the model weights, gets shared. FlashMLA, DeepEP and DeepGEMM together exposed a stack tuned for training and serving sparse [mixture-of-experts](https://aiwiki.ai/wiki/mixture_of_experts) models, and the Ascend release of September 2026 extended that pattern to a second hardware vendor. [2][14]

## What are FlashMLA's limitations?

FlashMLA is specialized, and that is its main constraint. The current version runs only on NVIDIA SM100 and SM103 GPUs and the Ascend 950 NPU, and only the sparse kernels exist for Ascend. It is written for the model DeepSeek is serving now: the September 2026 release dropped V3, V3.2 and V4.0, so anyone running those has to pin an older commit. Its decoding kernels accept only FP8 and FP4 caches, so a deployment that wants an unquantized cache cannot use them. The fused kernel, which the README recommends on CUDA, requires permuting two weight matrices ahead of time and pulls in TileLang, TileKernels and DeepGEMM for its test path. [1]

The performance figures are best read as a ceiling for a favorable configuration on one specific accelerator rather than a guarantee for any workload, and they are self-reported. And like any kernel that lives close to the hardware, FlashMLA carries tight coupling to specific toolchain versions and to vendor-specific features, which is the cost of running near the limits of the machine. [1][3]

## References

1. DeepSeek. "FlashMLA: Efficient Multi-head Latent Attention Kernels." GitHub repository and README. https://github.com/deepseek-ai/FlashMLA
2. DeepSeek. "Hello, DeepSeek Open Infra!" open-infra-index, GitHub. https://github.com/deepseek-ai/open-infra-index
3. DeepSeek. "A Deep-Dive Into the New Flash MLA Kernel." 22 April 2025. https://github.com/deepseek-ai/FlashMLA/blob/main/docs/20250422-new-kernel-deep-dive.md
4. DeepSeek-AI. "DeepSeek-V2: A Strong, Economical, and Efficient Mixture-of-Experts Language Model." arXiv:2405.04434, 2024. https://arxiv.org/abs/2405.04434
5. DeepSeek. "A Deep Dive Into The Flash MLA FP8 Decoding Kernel on Hopper." 29 September 2025. https://github.com/deepseek-ai/FlashMLA/blob/main/docs/20250929-hopper-fp8-sparse-deep-dive.md
6. DeepSeek-AI. "DeepSeek-V3 Technical Report." arXiv:2412.19437, 2024. https://arxiv.org/abs/2412.19437
7. Shah, J., Bikshandi, G., Zhang, Y., Thakkar, V., Ramani, P., Dao, T. "FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision." arXiv:2407.08608, 2024. https://arxiv.org/abs/2407.08608
8. NVIDIA. "CUTLASS: CUDA Templates for Linear Algebra Subroutines." GitHub. https://github.com/NVIDIA/cutlass
9. Shazeer, N. "Fast Transformer Decoding: One Write-Head is All You Need." arXiv:1911.02150, 2019. https://arxiv.org/abs/1911.02150
10. Zelenovic, S. "Enhancing DeepSeek models with MLA and FP8 optimizations in vLLM." Red Hat blog, 4 February 2025. https://www.redhat.com/en/blog/enhancing-deepseek-models-mla-and-fp8-optimizations-vllm
11. LMSYS Org. "SGLang v0.3 Release: 7x Faster DeepSeek MLA, 1.5x Faster torch.compile, Multi-Image/Video LLaVA-OneVision." 4 September 2024. https://www.lmsys.org/blog/2024-09-04-sglang-v0-3/
12. DeepSeek. "A Deep Dive Into the Ascend Sparse Attention Forward Kernel." 30 September 2026. https://github.com/deepseek-ai/FlashMLA/blob/main/docs/20260930-ascend-prefill-deep-dive.md
13. "Open-source release for Ascend 950 (2026.09.30)." Pull request 229, deepseek-ai/FlashMLA, merged 30 September 2026. https://github.com/deepseek-ai/FlashMLA/pull/229
14. Feng, C. "China's DeepSeek open-sources tools to help Huawei chips supplant Nvidia in AI." South China Morning Post, 30 September 2026. https://www.scmp.com/tech/tech-trends/article/3369301/chinas-deepseek-open-sources-tools-help-huawei-chips-supplant-nvidia-ai
15. IT Home. "DeepSeek open-sources infrastructure components for Huawei Ascend computing platform." 30 September 2026, republished by ifeng. https://tech.ifeng.com/c/8wpwu6MVfRV
16. DeepSeek. "FlashMLA setup.py." GitHub. https://github.com/deepseek-ai/FlashMLA/blob/main/setup.py
17. DeepSeek. "DeepSeek-V4.1-Flash." Hugging Face. https://huggingface.co/deepseek-ai/DeepSeek-V4.1-Flash
18. vLLM Project. "DeepSeek-V3.2-Exp in vLLM: Fine-Grained Sparse Attention in Action." 29 September 2025. https://blog.vllm.ai/2025/09/29/deepseek-v3-2.html