feat(pytorch): add mimo-v2-flash support - #4950
yidingcheng0206 wants to merge 22 commits into
Conversation
There was a problem hiding this comment.
Note
Copilot was unable to run its full agentic suite in this review.
Pull request overview
Adds MiMo-V2-Flash support across serving + PyTorch engine, including MiMo-specific tool parsing, MTP speculative decoding, and CUDA attention/graph improvements.
Changes:
- Introduces MiMo-V2-Flash target + MTP draft model implementations, config builder, and module registration.
- Extends speculative decoding pipeline (proposer depth transitions, DP meta handling, CUDA Graph warmup/capture safety).
- Improves attention backends/kernels (paged multi-token causal decode, FA3 capability gating, flash-attn asymmetric head tuning) and adds extensive tests/docs.
Reviewed changes
Copilot reviewed 51 out of 51 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
| tests/test_lmdeploy/serve/parsers/test_mimo_parser.py | Adds tests for MiMo tool-call parsing (complete + streaming + filtering). |
| tests/pytorch/test_distributed.py | Adds regression test for reduce_scatter output aliasing. |
| tests/pytorch/spec_decode/test_spec_agent.py | Updates spec-agent tests for new proposer depth-transition interface and DP token layouts. |
| tests/pytorch/spec_decode/test_cudagraph_strategy.py | Adds tests for CUDA graph buffer fill immutability and capture rollback semantics. |
| tests/pytorch/nn/test_linear_blocked_fp8.py | Adds tests for checkpoint-aware FP8 QKV sharding/scale behavior. |
| tests/pytorch/kernel/test_swa_state_ring.py | Adds CUDA kernel tests for state-ring flatten/scatter and attention equivalence. |
| tests/pytorch/kernel/test_paged_attention.py | Adds causal multi-token paged attention test + split-K neutrality test; refactors fixtures/classes. |
| tests/pytorch/kernel/test_mla_attention.py | Extends MLA builder test coverage for allow_fa3=False. |
| tests/pytorch/kernel/test_flash_attention.py | Adds test for asymmetric (K=192,V=128) flash attention head shape. |
| tests/pytorch/kernel/test_fa3_attention.py | Adds builder fallback test when FA3 wheel lacks asymmetric head instantiation. |
| lmdeploy/serve/parsers/tool_parser/qwen3coder_tool_parser.py | Adjusts whitespace handling so parser can preserve raw parameter text when needed. |
| lmdeploy/serve/parsers/tool_parser/mimo_tool_parser.py | Adds MiMo tool parser (HTML unescape + schema coercion preserving string whitespace). |
| lmdeploy/serve/parsers/tool_parser/init.py | Exposes MiMo tool parser in package exports. |
| lmdeploy/pytorch/weight_loader/model_weight_loader.py | Allows models to select checkpoint shards via select_weight_paths hook. |
| lmdeploy/pytorch/strategies/ar_spec/cudagraph.py | Adds MiMo MTP CUDA graph max-token sizing logic. |
| lmdeploy/pytorch/spec_decode/spec_agent.py | Uses proposer advance_draft_depth + proposer-provided DP token layout; improves warmup graphs. |
| lmdeploy/pytorch/spec_decode/proposers/mimo_mtp.py | Implements MiMo “same-position multi-depth” proposer and shared-embedding binding. |
| lmdeploy/pytorch/spec_decode/proposers/base.py | Introduces advance_draft_depth and DP token layout hook for draft-depth transitions. |
| lmdeploy/pytorch/spec_decode/proposers/init.py | Registers MiMo proposer. |
| lmdeploy/pytorch/nn/linear/blocked_fp8.py | Adds checkpoint-aware, block-aligned FP8 QKV sharding plan + logical/physical split_qkv. |
| lmdeploy/pytorch/nn/linear/init.py | Plumbs checkpoint-aware QKV sharding options into builder. |
| lmdeploy/pytorch/nn/attention.py | Exposes allow_fa3 and enable_paged_multi_token_decode in Attention build spec. |
| lmdeploy/pytorch/models/utils/cudagraph.py | Adds warmup spec API + capture-cache hook; avoids mutating caller-owned attn_metadata in fill. |
| lmdeploy/pytorch/models/module_map.py | Registers MiMo-V2-Flash target + MTP model entries. |
| lmdeploy/pytorch/models/mimo_v2_flash_mtp.py | Implements MiMo MTP draft model (3-depth predictor), loading, and CUDA graph behavior. |
| lmdeploy/pytorch/models/mimo_v2_flash.py | Implements MiMo-V2-Flash target model with hybrid full/SWA attention, FP8 loading, paged verification. |
| lmdeploy/pytorch/models/deepseek_mtp.py | Propagates spec_step_idx through generation inputs. |
| lmdeploy/pytorch/model_inputs.py | Adds spec_step_idx to ModelInputs/StepContext and advances it on step. |
| lmdeploy/pytorch/kernels/cuda/swa_state_ring.py | Adds Triton flatten/scatter kernels for SWA BF16 state rings. |
| lmdeploy/pytorch/kernels/cuda/pagedattention.py | Adds causal multi-token mask path for paged attention + fixes empty split-K workspace init. |
| lmdeploy/pytorch/kernels/cuda/flashattention.py | Fixes kernel meta selection for split-D (e.g., 192) attention heads. |
| lmdeploy/pytorch/engine/config_builder.py | Validates/constructs mimo_mtp spec config and MiMo-V2-Flash constraints. |
| lmdeploy/pytorch/distributed.py | Ensures reduce_scatter output is disjoint from input views. |
| lmdeploy/pytorch/configurations/mimo_v2_flash.py | Adds MiMo-V2-Flash model config builder + cache policy (paged SWA vs ring). |
| lmdeploy/pytorch/backends/cuda/graph_runner/runner.py | Adds model capability gating for FA3 requirement and capture rollback snapshots for mutable caches. |
| lmdeploy/pytorch/backends/cuda/graph_runner/full_graph.py | Threads supports_non_fa3_speculative_graph into graph meta for FA3 usage. |
| lmdeploy/pytorch/backends/cuda/attention/swa_state_ring.py | Adds SWA attention backend using state ring + flatten/scatter fallback for q>1. |
| lmdeploy/pytorch/backends/cuda/attention/fa3_capabilities.py | Adds runtime check for FA3 wheel head-dim template availability. |
| lmdeploy/pytorch/backends/cuda/attention/fa3.py | Advertises paged multi-token decode support on FA3 implementation. |
| lmdeploy/pytorch/backends/cuda/attention/default.py | Adds optional paged multi-token decode support and kernel-metadata override in forward. |
| lmdeploy/pytorch/backends/cuda/attention/init.py | Gates FA3 enabling on allow_fa3 + wheel capability; plumbs multi-token decode option. |
| lmdeploy/pytorch/backends/attention.py | Extends PagedAttentionBuildSpec with allow_fa3 and enable_paged_multi_token_decode. |
| lmdeploy/cli/utils.py | Adds mimo_mtp to CLI speculative algorithm choices. |
| docs/zh_cn/supported_models/supported_models.md | Documents MiMo-V2-Flash support and constraints (BF16 runtime, CUDA backend). |
| docs/zh_cn/llm/api_server_tools.md | Documents MiMo tool-call parser usage in api_server. |
| docs/zh_cn/advance/spec_decoding.md | Documents MiMo-V2-Flash MTP configuration and backend behavior. |
| docs/en/supported_models/supported_models.md | Documents MiMo-V2-Flash support and constraints (BF16 runtime, CUDA backend). |
| docs/en/llm/api_server_tools.md | Documents MiMo tool-call parser usage in api_server. |
| docs/en/advance/spec_decoding.md | Documents MiMo-V2-Flash MTP configuration and backend behavior. |
| README_zh-CN.md | Lists MiMo-V2-Flash as supported model. |
| README.md | Lists MiMo-V2-Flash as supported model. |
Suppressed comments (1)
tests/pytorch/kernel/test_fa3_attention.py:1
- This test uses
sysandSimpleNamespacebut does not import them in the file. Addimport sysandfrom types import SimpleNamespace(or equivalent) so the test doesn't fail withNameError.
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
884412c to
a6753cb
Compare
48b8148 to
875cd79
Compare
875cd79 to
d28f0c3
Compare
| stride_boffb, | ||
| kv_group_num: tl.constexpr, | ||
| seq_len: tl.constexpr, | ||
| causal_multi_token: tl.constexpr, |
There was a problem hiding this comment.
There are large change in this kernel, It could be risky to do such large change in this kernel, we better create a new kernel for swa.
| output = runner.capture(**kwargs) | ||
| finally: | ||
| if capture_state is not None: | ||
| capture_state.restore() |
There was a problem hiding this comment.
I am not sure if it is worth to add these restore logical. the cache might already been polluted during capture.
| logits_indices: torch.LongTensor | None = None | ||
| # Number of compact logprob rows emitted for each sequence. | ||
| seq_logit_length: torch.LongTensor | None = None | ||
| # Prediction-depth index for multi-layer MTP draft models. |
There was a problem hiding this comment.
Do we need to set/update these values in strategy/model_inputs.py?
Motivation
This PR adds support for Xiaomi MiMo-V2-Flash in LMDeploy's PyTorch backend. MiMo-V2-Flash uses hybrid full-attention and sliding-window-attention (SWA) layers and provides a three-layer MTP module for speculative decoding. This PR adds the required model implementation, attention and KV-cache support, MTP integration, and tool-call parsing.
Modification
LMDeploy and Hugging Face were evaluated with aligned settings. The LMDeploy values are averages over repeated runs.
Acceptance of MTP
The following results use TP8 + DP2 + EP8, attention TP4 per DP group, fixed 128-request ShareGPT workload, 2048 generated tokens, temperature 0,
ignore_eos=true, three MTP draft positions, and three repeated runs.Aggregate results
Per-position results
LMDeploy's native acceptance rate is
accepted_draft_tokens / drafted_tokensand excludes the bonus token. Mean accept length includes the bonus token. Position values are prefix acceptance probabilities; they must not be multiplied together. The strict 3/3 rate is the probability that all three draft positions are accepted.Performance
The fixed-output benchmark uses the same ShareGPT inputs, TP8 + DP2 + EP8 topology, attention TP4 per DP group, concurrency 128, and three repetitions per point.
MTP throughput and TTFT
MTP TPOT and ITL
LMDeploy MTP output throughput is 96.5%, 95.8%, and 98.5% of SGLang MTP for 2048, 4096, and 8192 output tokens, respectively. MTP TPOT remains within 3% of SGLang in these tests.
Runtime validation
Compatibility
No intentional backward-incompatible changes are introduced. MiMo-V2-Flash support targets the PyTorch backend. MTP requires the three-layer draft weights included in the checkpoint. FA3 is optional; unsupported asymmetric head-dimension builds automatically use Triton.
Validation
git diff --checkpass.Assistance
Assisted with Codex, reviewed manually.
Checklist