Repository navigation
[Feature][Model] Support GLM-5.3-Flash on Ascend 950 - #15127
Conversation
Summary of ChangesHello, I'm Gemini Code Assist1! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed! This pull request enables support for the GLM-5.3-Flash model on Ascend 950. It introduces necessary infrastructure for KDA and MLA-NoPE operations, addresses specific issues with attention prefill paths and RMSNorm configurations, and provides comprehensive documentation and testing configurations to support the new model architecture. Highlights
New Features🧠 You can now enable Memory (public preview) to help Gemini Code Assist learn from your team's feedback. This makes future code reviews more consistent and personalized to your project's style. Click here to enable Memory in your admin console. Using Gemini Code AssistThe full guide for Gemini Code Assist can be found on our documentation page, here are some quick tips. Invoking Gemini You can request assistance from Gemini at any point by creating a comment using either
Customization To customize the Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a Limitations & Feedback Gemini Code Assist may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counterproductive. You can react with 👍 and 👎 on @gemini-code-assist comments. If you're interested in giving your feedback about your experience with Gemini Code Assist for GitHub and other Google products, sign up here. Footnotes
|
|
👋 Hi! Thank you for contributing to the vLLM Ascend project. The following points will speed up your PR merge:
If CI fails, you can run linting and testing checks locally according Contributing and Testing. |
|
This pull request has conflicts, please resolve those before we can evaluate the pull request. |
There was a problem hiding this comment.
Code Review
Suggested PR Title:
[Attention][Feature] Add native FP8 serving support for GLM-5.3-Flash on Ascend 950Suggested PR Summary:
### What this PR does / why we need it?
This PR introduces native FP8 serving support for the multimodal Mixture-of-Experts model `GLM-5.3-Flash` on Ascend 950 NPUs. It integrates support for GLM-5.3-Flash's hybrid architecture, which combines Kimi Delta Attention (KDA) linear-attention layers with NoPE sparse MLA layers.
Key changes include:
- Rebinding upstream Flash Linear Attention (FLA) KDA entry points to NPU Triton implementations.
- Implementing host-side KDA gate calculations (`apply_kda_gate`) and handling `qk_rope_head_dim == 0` (NoPE) in MLA.
- Ensuring GLM-5.3-Flash kpool indexer and tail cache layers retain their custom specifications and bypass DeepSeek SFA rewrites.
- Updating speculative decoding to support `Glm5NextForConditionalGeneration` and resolve its MTP layers.
- Adding documentation, configuration templates, and unit tests.
**Review Feedback:**
- **`mla_v1.py`**: The `cache.is_contiguous()` check will raise a `RuntimeError` when `page_size_padded` is used because the KV cache is allocated as a non-contiguous tensor using `torch.as_strided`. Use advanced indexing to support both contiguous and non-contiguous caches.
- **`kda.py` & `gate.py`**: Ensure tensors (like `beta` and the output of `apply_kda_gate`) are cast back to their original dtypes after intermediate `float32` operations to prevent dtype mismatches in Triton/NPU kernels.
- **`llm_base_proposer.py`**: Use defensive `getattr` checks when accessing `layer_module.shared_head.head` to prevent potential `AttributeError` crashes.
### Does this PR introduce _any_ user-facing change?
Yes, it adds support for serving `GLM-5.3-Flash` on Ascend 950 NPUs. Users can now serve this model using `vllm serve` with `--tensor-parallel-size 8` and `--trust-remote-code`.
### How was this patch tested?
- Added unit tests for KDA gate operations (`test_kda_gate.py`).
- Added unit tests for GLM-5.3-Flash configuration and quantization mapping (`test_fp8_config.py`, `test_modelslim_config.py`).
- Added end-to-end configuration template (`GLM-5.3-Flash.yaml`).f5972b5 to
97f60d4
Compare
4e3aa22 to
9b6cf00
Compare
90b3038 to
9b6cf00
Compare
|
This pull request has conflicts, please resolve those before we can evaluate the pull request. |
cb18e14 to
ecf6b47
Compare
fixed |
|
This pull request has conflicts, please resolve those before we can evaluate the pull request. |
|
Rebased onto latest main — the conflict flag should clear. Also fixed the CI failure, which was a regression from this PR rather than a flake. Cause: this PR set Fix: restore upstream's False default and opt in explicitly only where the kpool layout needs it. AI assistance was used for the CI triage and this fix; I have reviewed every changed line. |
d4aa2f5 to
758f017
Compare
|
This pull request has conflicts, please resolve those before we can evaluate the pull request. |
Port the GLM-5.3-Flash architecture (hybrid KDA linear attention, NoPE sparse MLA, native FP8) into vllm_ascend/models/glm5next/ and register it through ModelRegistry, so serving needs only vLLM main plus this branch instead of an unmerged upstream model definition. Three helpers that upstream satisfies with CUDA kernels get Ascend equivalents: the KDA chunk and recurrent entry points, a torch fwht128_quant_fp8 against a cached normalized Hadamard matrix, and index_select / index_copy_ in place of the CUDA-only state gather and scatter. Defects fixed, each with the symptom that identifies it: - mHC dropped every layer's input RMSNorm on out-of-tree platforms, because forward_oot reaches forward_native while mhc_pre_torch takes no norm_weight / norm_eps. Symptom: fluent but incoherent text, no NaNs and plausible activation magnitudes. - MLA-NoPE left the latent output unfilled in the chunked-context prefill path. The rope cache and its gather output are both zero-width, so npu_gather_pa_kv_cache returned early and any prompt longer than max_num_batched_tokens attended over uninitialised KV. Symptom: prompts return only '!' and GSM8K drops to 0.040. - Rebinding a patched op on its defining module missed callers that had already done a from-import, so the GLM vision tower kept the CUDA fused Q/K RMSNorm kernel and crashed on tl.extra.cuda.gdc_wait. The patch now sweeps sys.modules, guarding each module so a frozen or read-only one cannot abort startup. - The MTP path assumed a DeepSeek-shaped, text-only target model. The target lm_head is now resolved through the nested language model that multimodal wrappers such as Glm5NextForConditionalGeneration use, and _draft_embed_accepts_mm treats a head as text-only when inspect.signature cannot introspect it. Config registration stays a patch because the architecture is downstream: _CONFIG_REGISTRY learns the three GLM-5.3-Flash config classes and is_deepseek_mla is widened past its hard-coded model_type whitelist. Both are documented in vllm_ascend/patch/__init__.py with a removal condition. Unit tests cover the chunked-gather fix, the KDA gate, the MTP lm_head lookup and the multimodal-kwargs fallback. Signed-off-by: yiminghub2024 <482890@qq.com> Co-authored-by: Cursor <cursoragent@cursor.com>
AscendMLAAttentionSpec defaulted indexes_kv_by_block_stride to True, which opted every MLA model on the v1 runner into padded-page unification. On GLM-5.2 + dspark (TP8) that let the SFA indexer spec and the main MLA spec share one KV cache tensor, so _allocate_kv_cache_tensors handed the MLA layer the indexer's single-tensor tuple and _reshape_kv_cache_tensors failed with "not enough values to unpack (expected 2, got 1)". Restore upstream's False default and opt in explicitly for the kpool models that need it. model_uses_kpool_indexer also replaces the duplicated index_kpool attribute probes in enable_sfa / model_uses_sfa_sparse. Signed-off-by: yiminghub2024 <482890@qq.com> Co-authored-by: Cursor <cursoragent@cursor.com>
758f017 to
f41bb29
Compare
|
/rerun Rerun (failed jobs only):
|
1 similar comment
|
/rerun Rerun (failed jobs only):
|
| @@ -64,7 +64,7 @@ def causal_conv1d_fn( | |||
| pad_slot_id: int = PAD_SLOT_ID, | |||
|
/rerun Rerun (failed jobs only):
|
1 similar comment
|
/rerun Rerun (failed jobs only):
|
) ### What this PR does / why we need it? Adapts vllm-ascend for GLM-5.3-Flash (hybrid KDA linear attention + NoPE sparse MLA, native FP8) on Ascend 950. The architecture itself now lives in this repository, under `vllm_ascend/models/glm5next/`, so the PR no longer depends on the unmerged upstream GLM-5.3-Flash PR (vllm-project/vllm#53906). Every `vllm` symbol the ported code imports — 139 of them — resolves against vLLM main. What the port required on the Ascend side is described in *Porting the architecture downstream* below. Three defects made the model emit fluent-looking but wrong output, or fail only in environments other than the author's, so they are worth calling out: 1. `MHCPreOp` / `MHCFusedPostPreOp` accept `norm_weight` / `norm_eps` in `forward_native`, but `mhc_pre_torch` has no such parameters. Out-of-tree platforms reach `forward_native` through `forward_oot`, so all 45 layers ran without their input RMSNorm. No NaNs, plausible activation magnitudes, incoherent text. 2. In the chunked-context prefill path, MLA-NoPE leaves both the rope cache and its gather output zero-width, and `npu_gather_pa_kv_cache` then returns without filling the latent output. Any prompt longer than `max_num_batched_tokens` therefore attended over uninitialised KV. GSM8K dropped to 0.040. 3. Replacing an op by rebinding its defining module does not reach callers that already did a from-import. The GLM vision tower is imported during model registration, before `adapt_patch` runs, so it kept the CUDA fused Q/K RMSNorm kernel and crashed on `tl.extra.cuda.gdc_wait`. The patch now sweeps `sys.modules`. Full list of defects this PR fixes, each with the symptom that identifies it: | Symptom | Root cause | Where | | --- | --- | --- | | Fluent but incoherent text; no NaNs, plausible activations | `mhc_pre_torch` has no `norm_weight` / `norm_eps`, so `forward_native` dropped every layer's input RMSNorm on out-of-tree platforms | `patch/worker/patch_triton.py` | | Prompts longer than `max_num_batched_tokens` return only `!`; GSM8K 0.040 | rope cache and its gather output are both zero-width, so `npu_gather_pa_kv_cache` leaves the latent output unfilled | `attention/mla_v1.py` | | Vision tower raises `module 'triton.language.extra.cuda' has no attribute 'gdc_wait'` | the module was imported during model registration, before the patch ran, so its from-import kept the CUDA kernel | `patch/worker/patch_triton.py` | | FIA rejects `input_layout BNSD_NBSD`; reshape of a 0-element tensor | `qk_rope_head_dim == 0` leaves the rope operands empty and FIA falls back to GQA | `attention/mla_v1.py`, `ops/mla.py`, `ops/rotary_embedding.py` | | `ValueError: too many values to unpack (expected 2)` while loading weights | CANN 9.1 `npu_dynamic_mx_quant` returns `[out, in // 64, 2]`, the MXFP8 method unpacks `[out, in // 32]` | `quantization/methods/w8a8/fp8_block.py` | | `AssertionError` on `state.is_cuda` in the KDA path | `gather_initial_states` / `scatter_states` are CUDA-only and their kernels reference `gdc_wait` | `patch/worker/patch_triton.py` | | `TypeError: unsupported operand type(s) for -: 'NoneType' and 'int'` at scheduler init | the hybrid coordinator receives `num_prefill_lookahead=None` while the non-hybrid path guards it | `patch/platform/patch_kv_cache_coordinator.py` | | `block_size must be divisible by hash_block_size` | the GLM kpool tail uses `block_size=index_kpool` and opts out of prefix caching | `patch/platform/patch_kv_cache_coordinator.py` | | `'NoneType' object has no attribute 'decode'` during ACL graph update | hybrid KDA+MLA models register every layer but only MLA layers contribute a captured op | `attention/mla_v1.py`, `worker/model_runner_v1.py` | | MTP fails to start (four separate interface mismatches) | the spec-decode path assumes a DeepSeek-shaped, text-only target model | `spec_decode/llm_base_proposer.py` | | ACL graph capture aborts with 107027 / stalls at decode-FULL | the 310P `causal_conv1d_update` fallback calls `.item()`, a host sync | `patch/worker/patch_triton.py` | ### Porting the architecture downstream The model definition, config classes, multimodal processor and MTP module are carried under `vllm_ascend/models/glm5next/` and registered through `ModelRegistry.register_model`. Three helpers that upstream satisfies with CUDA kernels have Ascend equivalents in `vllm_ascend/models/glm5next/ops/`: | Helper | Upstream | Here | | --- | --- | --- | | KDA chunk / recurrent | CUDA Triton in `third_party/flash_linear_attention` | the existing NPU Triton entry points in `vllm_ascend.ops.triton.kda`, imported directly | | `fwht128_quant_fp8` | fused Hadamard-128 rotation + block-128 ue8m0 FP8 quantization | torch against a cached normalized Hadamard matrix, keeping the bf16 materialization before quantizing; device-side only, so ACL-graph safe | | `gather_initial_states` / `scatter_states` | Triton kernels asserting `state.is_cuda` | plain `index_select` / `index_copy_`, no host sync | `causal_conv1d` routes to the Ascend kernels through a local wrapper that drops the CUDA-side cache-management kwargs once instead of at each call site. `KpoolTailSpec` / `KpoolTailManager` register through vLLM's `register_custom_kv_cache_specs` platform hook, which runs after the built-in specs, so the kpool tail layers keep their own spec without patching the registry. Config registration still needs a patch, because the architecture is downstream: `_CONFIG_REGISTRY` learns the three GLM-5.3-Flash config classes, and `is_deepseek_mla` is widened because upstream resolves it from a hard-coded `model_type` whitelist. Both are documented in `vllm_ascend/patch/__init__.py` with a removal condition. Two consequences worth calling out: - The monkey patches that rebound module attributes on `vllm.models.glm5next.nvidia.kda` are now unreachable and have been removed. The remaining patches in `patch_triton.py` target upstream modules that the ported code still goes through (mHC, the vision tower's fused Q/K RMSNorm, the mamba state ops, FLA's KDA entry points), so they stay. - The kpool sparse indexer keeps its `CustomOp` layer so the index-K and tail caches still build, but the scoring path raises on Ascend. Upstream implements it as fused CUDA kernels (DeepGEMM block-FP8 MQA logits, paged MQA logits, radix top-k over a device workspace) with no NPU equivalent yet. GLM-5.3-Flash only enables it when a checkpoint sets `index_topk`; the configuration validated here is dense NoPE MLA plus KDA. ### Does this PR introduce any user-facing change? No new flags. GLM-5.3-Flash requires `--block-size 512`, since the kpool indexer uses `index_kpool * 32`. Serving no longer requires a vLLM branch carrying #53906; vLLM main plus this branch is enough. ### How was this patch tested? Ascend 950PR x8, TP=8, native FP8 checkpoint, `--max-model-len 131072`, `--block-size 512`, ACLGraph enabled. Everything below was re-run in a **brand-new container**, with vllm-ascend installed from this branch as it stood before the port and vLLM pinned to a fixed commit, to make sure no local state from development was involved. ACLGraph captured cleanly (piecewise 4/4, decode FULL 4/4 in 20 s) and prefill reached 6399.8 tokens/s, matching the development environment. That run therefore resolved the model from a vLLM branch carrying #53906, on the same checkpoint and the same kernels. The port relocates the model definition and swaps the three helpers listed above, leaving the dense NoPE MLA + KDA execution path unchanged, so a re-run on the ported tree is the remaining thing to confirm. Accuracy, GSM8K 5-shot via lm-eval `local-chat-completions`, 250 questions, greedy: | Configuration | flexible-extract | strict-match | | --- | ---: | ---: | | before the chunked-gather fix, no MTP | 0.040 | 0.036 | | after the fix, no MTP | 0.896 | 0.860 | | after the fix, MTP `num_speculative_tokens=3` | 0.880 | 0.860 | | after the fix, MTP=5, clean container rebuild | 0.892 | 0.856 | The first row is the regression signature of defect 2. The last three agree within one standard error (±0.02). Correctness of the chunked-context path, as a single deterministic request: a 28282-token prompt asking for the last number in the text answers `1200` exactly. Before the fix the same prompt returned only `!` tokens. AIME2026: aime26 report table: ┌───────────────┬───────────┬────────────┬──────────┬───────┬─────────┐ │ Model │ Dataset │ Metric │ Subset │ Num │ Score │ ├───────────────┼───────────┼────────────┼──────────┼───────┼─────────┤ │ glm-5.3-flash │ AIME-2026 │ Accuracy ↑ │ default │ 30 │ 86.7% │ └───────────────┴───────────┴────────────┴──────────┴───────┴─────────┘ gpqa_diamond report table: ┌───────────────┬──────────────┬────────────┬──────────┬───────┬─────────┐ │ Model │ Dataset │ Metric │ Subset │ Num │ Score │ ├───────────────┼──────────────┼────────────┼──────────┼───────┼─────────┤ │ glm-5.3-flash │ GPQA-Diamond │ Accuracy ↑ │ default │ 198 │ 77.8% │ Both numbers are lower bounds, in the same direction and for the same reason. Generations that exhaust the 40960-token budget while still inside the reasoning phase emit no final answer and score as wrong: 4 of 30 on AIME 2026, 36 of 198 on GPQA-Diamond. Their reasoning traces are not degenerate — no repetition loops, still advancing new case analysis at the cut — so this is the model's reasoning length on hard items, not a generation defect. I kept the cap rather than raising it, since capping and scoring over-length generations as wrong is the standard convention; scoring only the completed subset would be selection-biased upward and is not reported here. Performance: - Single request, real prompt, greedy: 55.6 tokens/s, MTP mean acceptance length 2.1-3.1. - 4 concurrent real prompts: acceptance 53% greedy, 30-44% at temperature=1.0 / top_p=0.95. - `vllm bench serve`, random dataset, 64K input / 3000 output, 10 requests: aggregate prefill 6400 tokens/s, MTP acceptance 46.21%, acceptance length 2.39. Acceptance on this dataset is a sensitive indicator of chunked-prefill correctness: before the fix it was 0%. Unit test added: `tests/ut/attention/test_mla_nope_chunked_gather.py` covers defect 2 and asserts the rope-present path is unchanged. Known dependencies: `Glm5NextForConditionalGeneration` and `Glm5NextMTPModel` are registered from this repository, so neither a vLLM fork nor a cherry-pick of #53906 is needed. - ACLGraph needs `causal_conv1d_update_npu`, which is not in vllm-ascend main yet. The binding added here is guarded and falls back when the kernel is absent, but decode-FULL capture then stalls on a host sync, so the kernel is a hard requirement for ACLGraph on KDA models. - vLLM main: vllm-project/vllm@ba07e4a --------- Signed-off-by: yiminghub2024 <482890@qq.com> Co-authored-by: Cursor <cursoragent@cursor.com>
) ### What this PR does / why we need it? Adapts vllm-ascend for GLM-5.3-Flash (hybrid KDA linear attention + NoPE sparse MLA, native FP8) on Ascend 950. The architecture itself now lives in this repository, under `vllm_ascend/models/glm5next/`, so the PR no longer depends on the unmerged upstream GLM-5.3-Flash PR (vllm-project/vllm#53906). Every `vllm` symbol the ported code imports — 139 of them — resolves against vLLM main. What the port required on the Ascend side is described in *Porting the architecture downstream* below. Three defects made the model emit fluent-looking but wrong output, or fail only in environments other than the author's, so they are worth calling out: 1. `MHCPreOp` / `MHCFusedPostPreOp` accept `norm_weight` / `norm_eps` in `forward_native`, but `mhc_pre_torch` has no such parameters. Out-of-tree platforms reach `forward_native` through `forward_oot`, so all 45 layers ran without their input RMSNorm. No NaNs, plausible activation magnitudes, incoherent text. 2. In the chunked-context prefill path, MLA-NoPE leaves both the rope cache and its gather output zero-width, and `npu_gather_pa_kv_cache` then returns without filling the latent output. Any prompt longer than `max_num_batched_tokens` therefore attended over uninitialised KV. GSM8K dropped to 0.040. 3. Replacing an op by rebinding its defining module does not reach callers that already did a from-import. The GLM vision tower is imported during model registration, before `adapt_patch` runs, so it kept the CUDA fused Q/K RMSNorm kernel and crashed on `tl.extra.cuda.gdc_wait`. The patch now sweeps `sys.modules`. Full list of defects this PR fixes, each with the symptom that identifies it: | Symptom | Root cause | Where | | --- | --- | --- | | Fluent but incoherent text; no NaNs, plausible activations | `mhc_pre_torch` has no `norm_weight` / `norm_eps`, so `forward_native` dropped every layer's input RMSNorm on out-of-tree platforms | `patch/worker/patch_triton.py` | | Prompts longer than `max_num_batched_tokens` return only `!`; GSM8K 0.040 | rope cache and its gather output are both zero-width, so `npu_gather_pa_kv_cache` leaves the latent output unfilled | `attention/mla_v1.py` | | Vision tower raises `module 'triton.language.extra.cuda' has no attribute 'gdc_wait'` | the module was imported during model registration, before the patch ran, so its from-import kept the CUDA kernel | `patch/worker/patch_triton.py` | | FIA rejects `input_layout BNSD_NBSD`; reshape of a 0-element tensor | `qk_rope_head_dim == 0` leaves the rope operands empty and FIA falls back to GQA | `attention/mla_v1.py`, `ops/mla.py`, `ops/rotary_embedding.py` | | `ValueError: too many values to unpack (expected 2)` while loading weights | CANN 9.1 `npu_dynamic_mx_quant` returns `[out, in // 64, 2]`, the MXFP8 method unpacks `[out, in // 32]` | `quantization/methods/w8a8/fp8_block.py` | | `AssertionError` on `state.is_cuda` in the KDA path | `gather_initial_states` / `scatter_states` are CUDA-only and their kernels reference `gdc_wait` | `patch/worker/patch_triton.py` | | `TypeError: unsupported operand type(s) for -: 'NoneType' and 'int'` at scheduler init | the hybrid coordinator receives `num_prefill_lookahead=None` while the non-hybrid path guards it | `patch/platform/patch_kv_cache_coordinator.py` | | `block_size must be divisible by hash_block_size` | the GLM kpool tail uses `block_size=index_kpool` and opts out of prefix caching | `patch/platform/patch_kv_cache_coordinator.py` | | `'NoneType' object has no attribute 'decode'` during ACL graph update | hybrid KDA+MLA models register every layer but only MLA layers contribute a captured op | `attention/mla_v1.py`, `worker/model_runner_v1.py` | | MTP fails to start (four separate interface mismatches) | the spec-decode path assumes a DeepSeek-shaped, text-only target model | `spec_decode/llm_base_proposer.py` | | ACL graph capture aborts with 107027 / stalls at decode-FULL | the 310P `causal_conv1d_update` fallback calls `.item()`, a host sync | `patch/worker/patch_triton.py` | ### Porting the architecture downstream The model definition, config classes, multimodal processor and MTP module are carried under `vllm_ascend/models/glm5next/` and registered through `ModelRegistry.register_model`. Three helpers that upstream satisfies with CUDA kernels have Ascend equivalents in `vllm_ascend/models/glm5next/ops/`: | Helper | Upstream | Here | | --- | --- | --- | | KDA chunk / recurrent | CUDA Triton in `third_party/flash_linear_attention` | the existing NPU Triton entry points in `vllm_ascend.ops.triton.kda`, imported directly | | `fwht128_quant_fp8` | fused Hadamard-128 rotation + block-128 ue8m0 FP8 quantization | torch against a cached normalized Hadamard matrix, keeping the bf16 materialization before quantizing; device-side only, so ACL-graph safe | | `gather_initial_states` / `scatter_states` | Triton kernels asserting `state.is_cuda` | plain `index_select` / `index_copy_`, no host sync | `causal_conv1d` routes to the Ascend kernels through a local wrapper that drops the CUDA-side cache-management kwargs once instead of at each call site. `KpoolTailSpec` / `KpoolTailManager` register through vLLM's `register_custom_kv_cache_specs` platform hook, which runs after the built-in specs, so the kpool tail layers keep their own spec without patching the registry. Config registration still needs a patch, because the architecture is downstream: `_CONFIG_REGISTRY` learns the three GLM-5.3-Flash config classes, and `is_deepseek_mla` is widened because upstream resolves it from a hard-coded `model_type` whitelist. Both are documented in `vllm_ascend/patch/__init__.py` with a removal condition. Two consequences worth calling out: - The monkey patches that rebound module attributes on `vllm.models.glm5next.nvidia.kda` are now unreachable and have been removed. The remaining patches in `patch_triton.py` target upstream modules that the ported code still goes through (mHC, the vision tower's fused Q/K RMSNorm, the mamba state ops, FLA's KDA entry points), so they stay. - The kpool sparse indexer keeps its `CustomOp` layer so the index-K and tail caches still build, but the scoring path raises on Ascend. Upstream implements it as fused CUDA kernels (DeepGEMM block-FP8 MQA logits, paged MQA logits, radix top-k over a device workspace) with no NPU equivalent yet. GLM-5.3-Flash only enables it when a checkpoint sets `index_topk`; the configuration validated here is dense NoPE MLA plus KDA. ### Does this PR introduce any user-facing change? No new flags. GLM-5.3-Flash requires `--block-size 512`, since the kpool indexer uses `index_kpool * 32`. Serving no longer requires a vLLM branch carrying #53906; vLLM main plus this branch is enough. ### How was this patch tested? Ascend 950PR x8, TP=8, native FP8 checkpoint, `--max-model-len 131072`, `--block-size 512`, ACLGraph enabled. Everything below was re-run in a **brand-new container**, with vllm-ascend installed from this branch as it stood before the port and vLLM pinned to a fixed commit, to make sure no local state from development was involved. ACLGraph captured cleanly (piecewise 4/4, decode FULL 4/4 in 20 s) and prefill reached 6399.8 tokens/s, matching the development environment. That run therefore resolved the model from a vLLM branch carrying #53906, on the same checkpoint and the same kernels. The port relocates the model definition and swaps the three helpers listed above, leaving the dense NoPE MLA + KDA execution path unchanged, so a re-run on the ported tree is the remaining thing to confirm. Accuracy, GSM8K 5-shot via lm-eval `local-chat-completions`, 250 questions, greedy: | Configuration | flexible-extract | strict-match | | --- | ---: | ---: | | before the chunked-gather fix, no MTP | 0.040 | 0.036 | | after the fix, no MTP | 0.896 | 0.860 | | after the fix, MTP `num_speculative_tokens=3` | 0.880 | 0.860 | | after the fix, MTP=5, clean container rebuild | 0.892 | 0.856 | The first row is the regression signature of defect 2. The last three agree within one standard error (±0.02). Correctness of the chunked-context path, as a single deterministic request: a 28282-token prompt asking for the last number in the text answers `1200` exactly. Before the fix the same prompt returned only `!` tokens. AIME2026: aime26 report table: ┌───────────────┬───────────┬────────────┬──────────┬───────┬─────────┐ │ Model │ Dataset │ Metric │ Subset │ Num │ Score │ ├───────────────┼───────────┼────────────┼──────────┼───────┼─────────┤ │ glm-5.3-flash │ AIME-2026 │ Accuracy ↑ │ default │ 30 │ 86.7% │ └───────────────┴───────────┴────────────┴──────────┴───────┴─────────┘ gpqa_diamond report table: ┌───────────────┬──────────────┬────────────┬──────────┬───────┬─────────┐ │ Model │ Dataset │ Metric │ Subset │ Num │ Score │ ├───────────────┼──────────────┼────────────┼──────────┼───────┼─────────┤ │ glm-5.3-flash │ GPQA-Diamond │ Accuracy ↑ │ default │ 198 │ 77.8% │ Both numbers are lower bounds, in the same direction and for the same reason. Generations that exhaust the 40960-token budget while still inside the reasoning phase emit no final answer and score as wrong: 4 of 30 on AIME 2026, 36 of 198 on GPQA-Diamond. Their reasoning traces are not degenerate — no repetition loops, still advancing new case analysis at the cut — so this is the model's reasoning length on hard items, not a generation defect. I kept the cap rather than raising it, since capping and scoring over-length generations as wrong is the standard convention; scoring only the completed subset would be selection-biased upward and is not reported here. Performance: - Single request, real prompt, greedy: 55.6 tokens/s, MTP mean acceptance length 2.1-3.1. - 4 concurrent real prompts: acceptance 53% greedy, 30-44% at temperature=1.0 / top_p=0.95. - `vllm bench serve`, random dataset, 64K input / 3000 output, 10 requests: aggregate prefill 6400 tokens/s, MTP acceptance 46.21%, acceptance length 2.39. Acceptance on this dataset is a sensitive indicator of chunked-prefill correctness: before the fix it was 0%. Unit test added: `tests/ut/attention/test_mla_nope_chunked_gather.py` covers defect 2 and asserts the rope-present path is unchanged. Known dependencies: `Glm5NextForConditionalGeneration` and `Glm5NextMTPModel` are registered from this repository, so neither a vLLM fork nor a cherry-pick of #53906 is needed. - ACLGraph needs `causal_conv1d_update_npu`, which is not in vllm-ascend main yet. The binding added here is guarded and falls back when the kernel is absent, but decode-FULL capture then stalls on a host sync, so the kernel is a hard requirement for ACLGraph on KDA models. - vLLM main: vllm-project/vllm@ba07e4a --------- Signed-off-by: yiminghub2024 <482890@qq.com> Co-authored-by: Cursor <cursoragent@cursor.com>
…oject#14620 vllm-project#14620 (0825) removed the Ascend Triton causal_conv1d_update_npu as "obsolete"; vllm-project#15127 (0903) then re-wired patch_triton.py to import it for the GLM-5.3-Flash KDA layers, leaving every 0.29-line deployment binding the PyTorch fallback instead (patch_triton.py:326 ImportError -> :332 warning -> host-syncing fallback on every decode step). Kernel body is the upstream tiled implementation from f0a9389~1 (588-line tree), kept line-identical. Wrapper is re-bound to the vllm 0.28/0.29 call contract: - stride-rebinding writes conv_state updates in place instead of the upstream transpose().contiguous() copies (which silently dropped state updates); - null_block_id defaults to the NULL_BLOCK_ID sentinel (int, not None) so the non-spec decode path (validate_data=True, no explicit null_block_id, e.g. kda.py:980) passes the assertion, matching the upstream CUDA/Triton wrapper; Parity vs the vllm upstream Triton kernel on NPU (decoder-0, Triton 3.2.0 / torch 2.10 / CANN 9.1.0), all bit-exact on out+state: 2D non-spec dim512/w4 and dim768/w2; spec varlen two-step sliding window (MAL=4, accepted 4/2/1/3 -> 2/4/1/4); spec 3D; non-spec 3D; null-block sentinel; fp32 observed group only ~1e-3 (fp16 midpoint, in-tolerance). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…oject#14620 vllm-project#14620 (0825) removed the Ascend Triton causal_conv1d_update_npu as "obsolete"; vllm-project#15127 (0903) then re-wired patch_triton.py to import it for the GLM-5.3-Flash KDA layers, leaving every 0.29-line deployment binding the PyTorch fallback instead (patch_triton.py:326 ImportError -> :332 warning -> host-syncing fallback on every decode step). Kernel body is the upstream tiled implementation from f0a9389~1 (588-line tree), kept line-identical. Wrapper is re-bound to the vllm 0.28/0.29 call contract: - stride-rebinding writes conv_state updates in place instead of the upstream transpose().contiguous() copies (which silently dropped state updates); - null_block_id defaults to the NULL_BLOCK_ID sentinel (int, not None) so the non-spec decode path (validate_data=True, no explicit null_block_id, e.g. kda.py:980) passes the assertion, matching the upstream CUDA/Triton wrapper; Parity vs the vllm upstream Triton kernel on NPU (decoder-0, Triton 3.2.0 / torch 2.10 / CANN 9.1.0), all bit-exact on out+state: 2D non-spec dim512/w4 and dim768/w2; spec varlen two-step sliding window (MAL=4, accepted 4/2/1/3 -> 2/4/1/4); spec 3D; non-spec 3D; null-block sentinel; fp32 observed group only ~1e-3 (fp16 midpoint, in-tolerance). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
What this PR does / why we need it?
Adapts vllm-ascend for GLM-5.3-Flash (hybrid KDA linear attention + NoPE
sparse MLA, native FP8) on Ascend 950.
The architecture itself now lives in this repository, under
vllm_ascend/models/glm5next/, so the PR no longer depends on the unmergedupstream GLM-5.3-Flash PR (vllm-project/vllm#53906). Every
vllmsymbol theported code imports — 139 of them — resolves against vLLM main. What the port
required on the Ascend side is described in Porting the architecture
downstream below.
Three defects made the model emit fluent-looking but wrong output, or fail
only in environments other than the author's, so they are worth calling out:
MHCPreOp/MHCFusedPostPreOpacceptnorm_weight/norm_epsinforward_native, butmhc_pre_torchhas no such parameters. Out-of-treeplatforms reach
forward_nativethroughforward_oot, so all 45 layersran without their input RMSNorm. No NaNs, plausible activation magnitudes,
incoherent text.
and its gather output zero-width, and
npu_gather_pa_kv_cachethen returnswithout filling the latent output. Any prompt longer than
max_num_batched_tokenstherefore attended over uninitialised KV. GSM8Kdropped to 0.040.
that already did a from-import. The GLM vision tower is imported during
model registration, before
adapt_patchruns, so it kept the CUDA fusedQ/K RMSNorm kernel and crashed on
tl.extra.cuda.gdc_wait. The patch nowsweeps
sys.modules.Full list of defects this PR fixes, each with the symptom that identifies it:
mhc_pre_torchhas nonorm_weight/norm_eps, soforward_nativedropped every layer's input RMSNorm on out-of-tree platformspatch/worker/patch_triton.pymax_num_batched_tokensreturn only!; GSM8K 0.040npu_gather_pa_kv_cacheleaves the latent output unfilledattention/mla_v1.pymodule 'triton.language.extra.cuda' has no attribute 'gdc_wait'patch/worker/patch_triton.pyinput_layout BNSD_NBSD; reshape of a 0-element tensorqk_rope_head_dim == 0leaves the rope operands empty and FIA falls back to GQAattention/mla_v1.py,ops/mla.py,ops/rotary_embedding.pyValueError: too many values to unpack (expected 2)while loading weightsnpu_dynamic_mx_quantreturns[out, in // 64, 2], the MXFP8 method unpacks[out, in // 32]quantization/methods/w8a8/fp8_block.pyAssertionErroronstate.is_cudain the KDA pathgather_initial_states/scatter_statesare CUDA-only and their kernels referencegdc_waitpatch/worker/patch_triton.pyTypeError: unsupported operand type(s) for -: 'NoneType' and 'int'at scheduler initnum_prefill_lookahead=Nonewhile the non-hybrid path guards itpatch/platform/patch_kv_cache_coordinator.pyblock_size must be divisible by hash_block_sizeblock_size=index_kpooland opts out of prefix cachingpatch/platform/patch_kv_cache_coordinator.py'NoneType' object has no attribute 'decode'during ACL graph updateattention/mla_v1.py,worker/model_runner_v1.pyspec_decode/llm_base_proposer.pycausal_conv1d_updatefallback calls.item(), a host syncpatch/worker/patch_triton.pyPorting the architecture downstream
The model definition, config classes, multimodal processor and MTP module are
carried under
vllm_ascend/models/glm5next/and registered throughModelRegistry.register_model. Three helpers that upstream satisfies with CUDAkernels have Ascend equivalents in
vllm_ascend/models/glm5next/ops/:third_party/flash_linear_attentionvllm_ascend.ops.triton.kda, imported directlyfwht128_quant_fp8gather_initial_states/scatter_statesstate.is_cudaindex_select/index_copy_, no host synccausal_conv1droutes to the Ascend kernels through a local wrapper that dropsthe CUDA-side cache-management kwargs once instead of at each call site.
KpoolTailSpec/KpoolTailManagerregister through vLLM'sregister_custom_kv_cache_specsplatform hook, which runs after the built-inspecs, so the kpool tail layers keep their own spec without patching the
registry.
Config registration still needs a patch, because the architecture is
downstream:
_CONFIG_REGISTRYlearns the three GLM-5.3-Flash config classes,and
is_deepseek_mlais widened because upstream resolves it from a hard-codedmodel_typewhitelist. Both are documented invllm_ascend/patch/__init__.pywith a removal condition.
Two consequences worth calling out:
vllm.models.glm5next.nvidia.kdaare now unreachable and have been removed.The remaining patches in
patch_triton.pytarget upstream modules that theported code still goes through (mHC, the vision tower's fused Q/K RMSNorm,
the mamba state ops, FLA's KDA entry points), so they stay.
CustomOplayer so the index-K and tailcaches still build, but the scoring path raises on Ascend. Upstream
implements it as fused CUDA kernels (DeepGEMM block-FP8 MQA logits, paged MQA
logits, radix top-k over a device workspace) with no NPU equivalent yet.
GLM-5.3-Flash only enables it when a checkpoint sets
index_topk; theconfiguration validated here is dense NoPE MLA plus KDA.
Does this PR introduce any user-facing change?
No new flags. GLM-5.3-Flash requires
--block-size 512, since the kpoolindexer uses
index_kpool * 32. Serving no longer requires a vLLM branchcarrying #53906; vLLM main plus this branch is enough.
How was this patch tested?
Ascend 950PR x8, TP=8, native FP8 checkpoint,
--max-model-len 131072,--block-size 512, ACLGraph enabled.Everything below was re-run in a brand-new container, with vllm-ascend
installed from this branch as it stood before the port and vLLM pinned to a
fixed commit, to make sure no local state from development was involved.
ACLGraph captured cleanly (piecewise 4/4, decode FULL 4/4 in 20 s) and prefill
reached 6399.8 tokens/s, matching the development environment.
That run therefore resolved the model from a vLLM branch carrying #53906, on
the same checkpoint and the same kernels. The port relocates the model
definition and swaps the three helpers listed above, leaving the dense NoPE
MLA + KDA execution path unchanged, so a re-run on the ported tree is the
remaining thing to confirm.
Accuracy, GSM8K 5-shot via lm-eval
local-chat-completions, 250 questions,greedy:
num_speculative_tokens=3The first row is the regression signature of defect 2. The last three agree
within one standard error (±0.02).
Correctness of the chunked-context path, as a single deterministic request:
a 28282-token prompt asking for the last number in the text answers
1200exactly. Before the fix the same prompt returned only
!tokens.AIME2026:
aime26 report table:
┌───────────────┬───────────┬────────────┬──────────┬───────┬─────────┐
│ Model │ Dataset │ Metric │ Subset │ Num │ Score │
├───────────────┼───────────┼────────────┼──────────┼───────┼─────────┤
│ glm-5.3-flash │ AIME-2026 │ Accuracy ↑ │ default │ 30 │ 86.7% │
└───────────────┴───────────┴────────────┴──────────┴───────┴─────────┘
gpqa_diamond report table:
┌───────────────┬──────────────┬────────────┬──────────┬───────┬─────────┐
│ Model │ Dataset │ Metric │ Subset │ Num │ Score │
├───────────────┼──────────────┼────────────┼──────────┼───────┼─────────┤
│ glm-5.3-flash │ GPQA-Diamond │ Accuracy ↑ │ default │ 198 │ 77.8% │
Both numbers are lower bounds, in the same direction and for the same reason.
Generations that exhaust the 40960-token budget while still inside the reasoning
phase emit no final answer and score as wrong: 4 of 30 on AIME 2026, 36 of 198 on
GPQA-Diamond. Their reasoning traces are not degenerate — no repetition loops,
still advancing new case analysis at the cut — so this is the model's reasoning
length on hard items, not a generation defect. I kept the cap rather than raising
it, since capping and scoring over-length generations as wrong is the standard
convention; scoring only the completed subset would be selection-biased upward and
is not reported here.
Performance:
length 2.1-3.1.
/ top_p=0.95.
vllm bench serve, random dataset, 64K input / 3000 output, 10 requests:aggregate prefill 6400 tokens/s, MTP acceptance 46.21%, acceptance length
2.39. Acceptance on this dataset is a sensitive indicator of chunked-prefill
correctness: before the fix it was 0%.
Unit test added:
tests/ut/attention/test_mla_nope_chunked_gather.pycoversdefect 2 and asserts the rope-present path is unchanged.
Known dependencies:
Glm5NextForConditionalGenerationandGlm5NextMTPModelare registered fromthis repository, so neither a vLLM fork nor a cherry-pick of #53906 is
needed.
ACLGraph needs
causal_conv1d_update_npu, which is not in vllm-ascend mainyet. The binding added here is guarded and falls back when the kernel is
absent, but decode-FULL capture then stalls on a host sync, so the kernel is
a hard requirement for ACLGraph on KDA models.
vLLM main: vllm-project/vllm@ba07e4a