[CI]Main2Main 0807 - #13477
[CI]Main2Main 0807#13477
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 performs a routine synchronization of the local vllm dependency tracking file with the upstream vllm main branch. This ensures that the CI environment is aligned with the latest developments and fixes from the upstream repository. 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
|
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
There was a problem hiding this comment.
Code Review
This pull request updates the verified commit hash for the vllm-main dependency in .github/vllm-main-verified.commit to 72cd5424da80a4a9caa3f42fd65bc0b94e61cbf0. There are no review comments, so I have no feedback to provide.
Suggested PR Title:
[CI][Misc] Update vllm-main-verified commit hashSuggested PR Summary:
### What this PR does / why we need it?
This pull request updates the verified commit hash for the `vllm-main` dependency in `.github/vllm-main-verified.commit` to `72cd5424da80a4a9caa3f42fd65bc0b94e61cbf0`.
### Does this PR introduce _any_ user-facing change?
No.
### How was this patch tested?
No tests were added as this is a simple dependency commit hash update.
|
👋 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. |
8d04145 to
c1d02a4
Compare
fd22101 to
3da8ea9
Compare
3da8ea9 to
7ffba89
Compare
7ffba89 to
b9fe5e6
Compare
b9fe5e6 to
0e1c88d
Compare
0e1c88d to
5ec06ac
Compare
5ec06ac to
f1fd765
Compare
0e64c1d to
b346687
Compare
b346687 to
7f60753
Compare
| # main branch renamed FusedMoE -> FusedMoEFactory; v0.26.0 keeps FusedMoE. | ||
| # The FusedMoEFactory import is kept for re-export compatibility on v0.26.0 only. | ||
| if vllm_version_is("0.26.0"): | ||
| from vllm.model_executor.layers.fused_moe.layer import FusedMoE # noqa: F401 |
There was a problem hiding this comment.
we don't need it. Please delete it.
|
/rerun Failed:
|
| if self.routed_experts_initialized: | ||
| self.routed_experts_capturer.clear_buffer() | ||
| if vllm_version_is("0.26.0"): | ||
| self.routed_experts_capturer.clear_buffer() |
There was a problem hiding this comment.
This change is incorrect, it shoud be:
if vllm_version_is("0.26.0"):
if self.vllm_config.model_config.enable_return_routed_experts:
if self.routed_experts_initialized:
self.routed_experts_capturer.clear_buffer()
|
|
||
|
|
||
| def _patched_fused_input_norm_forward(self, grid_thw, visual_dtype): | ||
| if self.is_identity: |
There was a problem hiding this comment.
Please add patch descriptions to vllm_ascend/patch/__init__.py
| state_conv_widths_ptr, # conv width for conv states (0 for temporal) | ||
| state_group_indices_ptr, # maps state_idx to group index in block table | ||
| # DS conv row metadata. Zero keeps the single-region copy path. | ||
| state_dim_row_count_ptr, # int32: per-block dim row count for DS conv |
There was a problem hiding this comment.
Please have the model owner confirm that the modifications are correct and there is no performance degradation.
|
/rerun Rerun:
|
|
This pull request has conflicts, please resolve those before we can evaluate the pull request. |
|
/rerun [Bot]: rerun completed. Rerun:
|
|
/rerun Rerun:
|
|
/rerun Failed:
|
| @@ -1 +1 @@ | |||
| # Adapt from https://github.com/vllm-project/vllm/blob/main/vllm/v1/worker/mamba_utils.py | |||
There was a problem hiding this comment.
好的，后面补充一下该用例和MD覆盖到涉及的场景
Update vllm-main-verified.commit to track upstream vllm main commit
8d9b52f7c2514490bdadfd5eb0c931e58625df2e.
Adaptations for upstream main changes between 2e09247c and 8d9b52f7c:
- FusedMoE → FusedMoEFactory rename (vllm#44941): version-gated patch in
patch_fused_moe.py, updated imports/usage in deepseek_v4.py,
minimax_m3.py, fused_moe.py, routed_experts.py, and test monkeypatch
- Removed calculate_kv_scales (vllm#49389): version-gated in
model_runner_v1.py
- Removed Q/K/V_SCALE_CONSTANT env var registrations (vllm#49389):
remove unused dead-code from attention layer.py. Module-level typed
constants still exist in envs.py.
- RoutedExpertsCapturer.clear_buffer() removed (vllm#50721): version-gated
with vllm_version_is("0.26.0") and merged inner conditions
- FusedInputNorm eps=0.0 (vllm#50411): patch FusedInputNorm.forward to
replace eps=0.0 with eps=1e-5 for Ascend torch_npu compatibility.
Guarded with contextlib.suppress(ImportError) for release wheels.
- Triton postprocess_mamba_fused_kernel signature changed (vllm#50432):
add CONV_STATE_DIM_FIRST, HAS_IDX_MAPPING, PRECOMPUTED_NEW_COMPUTED
parameters to Ascend's custom kernel
- Mamba prefix cache enabled by default (vllm#50991): enable
chunked_prefill in extract_hidden_states E2E test
- DeepSeek V4 reasoning effort mapping (vllm#50580): version-gated
thinking test expectations
- HunyuanVL placeholder format changed (vllm#49691): version-gated in
conftest.py
- Remove unused maybe_calc_kv_scales mock from spec_decode test
- Remove stale FusedMoE re-export from fused_moe.py
Signed-off-by: liaoqidan <1107297340@qq.com>
The FusedMoEFactory adaptation in 93bc205 introduced an unused _DefaultAscendMoERunner/_DefaultAscendRoutedExperts block whose type[RoutedExperts] annotation references an undefined name, failing ruff F821 in pre-commit. The block was never referenced; 310P runner selection already happens via REGISTERED_ASCEND_OPS in vllm_ascend/utils.py. Signed-off-by: liaoqidan <1107297340@qq.com>
Clarify that upstream vLLM uses PyTorch 2.13.0 (eps>0 training, eps>=0 inference) while vllm-ascend bundles PyTorch 2.10.0 (eps>0 always), so upstream passing eps=0.0 works there but fails on Ascend. The patch can be removed once bundled PyTorch >= 2.13.0. Signed-off-by: liaoqidan <1107297340@qq.com>
Sync .github/vllm-main-verified.commit to the latest vLLM main HEAD. Signed-off-by: liaoqidan <1107297340@qq.com>
LoganJane
left a comment
LoganJane
left a comment
There was a problem hiding this comment.
Requesting changes for five P1 correctness and compatibility issues in vllm_ascend/ops/triton/mamba/postprocess.py. Details are provided inline. This review is intentionally limited to P1 findings.
| if not PRECOMPUTED_NEW_COMPUTED: | ||
| num_tokens_running_state = num_computed + num_scheduled - num_draft | ||
| else: | ||
| num_tokens_running_state = (new_num_computed // block_size) * block_size |
There was a problem hiding this comment.
[P1] Restore the MRv2 running-state formula.
Under PRECOMPUTED_NEW_COMPUTED, this assigns num_tokens_running_state to the same aligned value as aligned_new_computed. Consequently, needs_copy is always true and accept_token_bias is always zero. For example, with block_size=16, new_num_computed=32, and num_accepted=3, the correct values are num_tokens_running_state=30 and accept_token_bias=2, while this code produces 32 and 0. The upstream contract is:
new_num_computed = tl.load(num_computed_tokens_ptr + req_idx)
num_tokens_running_state = new_num_computed - num_accepted + 1Please use that formula; otherwise MRv2 align can skip required state shifts or copy at the wrong boundary.
| num_scheduled = tl.load(num_scheduled_tokens_ptr + req_idx) | ||
| num_computed = tl.load(num_computed_tokens_ptr + req_idx) | ||
| num_draft = tl.load(num_draft_tokens_ptr + req_idx) | ||
| if PRECOMPUTED_NEW_COMPUTED: |
There was a problem hiding this comment.
[P1] Do not load num_scheduled_tokens_ptr on the precomputed path.
run_fused_postprocess_align() passes None for both num_scheduled_tokens_ptr and num_draft_tokens_ptr when PRECOMPUTED_NEW_COMPUTED=True, but line 80 performs tl.load(num_scheduled_tokens_ptr + req_idx) unconditionally before this branch. This can fail during Triton specialization/compilation by doing pointer arithmetic and a load through None; relying on dead-code elimination is unsafe. Move the num_scheduled, num_computed, and num_draft loads entirely into the else branch, matching upstream.
| state_idx = tl.program_id(1) | ||
|
|
||
| if HAS_IDX_MAPPING: | ||
| req_idx = tl.load(idx_mapping_ptr + batch_idx).to(tl.int32) |
There was a problem hiding this comment.
[P1] Validate the batch index, not the mapped request slot, against num_reqs.
num_reqs is the number of active batch rows, whereas req_idx is a request-state slot and may be sparse/non-contiguous. For example, num_reqs=2 with idx_mapping=[5, 1] is valid, but the current req_idx >= num_reqs check drops slot 5 entirely. It also fails to reject the -1 skip sentinel and can read decision buffers at a negative offset. The control flow should be:
if batch_idx >= num_reqs:
return
if HAS_IDX_MAPPING:
req_idx = tl.load(idx_mapping_ptr + batch_idx)
if req_idx < 0:
return
else:
req_idx = batch_idxA positive request-slot bound would require the actual buffer capacity, not num_reqs.
| dim_row_count = tl.load(state_dim_row_count_ptr + state_idx) | ||
| dim_row_stride = tl.load(state_dim_row_stride_ptr + state_idx) | ||
| num_rows_to_copy = (conv_width - accept_token_bias).to(tl.int64) | ||
| copy_size = dim_row_count * dim_row_stride |
There was a problem hiding this comment.
[P1] Fix the DS convolution row copy dimensions and advance the row address.
For DS layout (state[block, dim, state_len]), the kernel must iterate dim_row_count rows and copy (conv_width - accept_token_bias) * state_elem_size bytes per row, advancing source and destination by dim_row_stride for each row. This code swaps those quantities: it uses conv_width - bias as the loop count and dim_row_count * dim_row_stride as the per-loop copy size. The loop at lines 192-196 then never applies a row offset, so it repeats the same oversized copy.
For dim=4, state_len=3, fp16, and bias=1, this copies 24 bytes from source offset 2, reading 2 bytes past a 24-byte block, and repeats it twice. Please use num_loops = dim_row_count, copy_size = (conv_width - bias) * state_elem_size, and add row * dim_row_stride to both pointers while keeping the pointer-type cast hoisted.
| CONV_STATE_DIM_FIRST: tl.constexpr, | ||
| # HAS_IDX_MAPPING: when True, program_id(0) is a batch index resolved to a | ||
| # req-state slot via idx_mapping_ptr (V2). When False, it is the req index. | ||
| HAS_IDX_MAPPING: tl.constexpr = False, |
There was a problem hiding this comment.
[P1] Preserve the vLLM v0.26.0 output-buffer contract.
This Ascend kernel is installed for both main and v0.26.0, but their MRv2 callers have different output contracts. Main (after vLLM #50432) passes an accepted-token snapshot as input and a non-null output buffer. v0.26.0 passes None as num_accepted_tokens_out_ptr and expects the HAS_IDX_MAPPING path to update num_accepted_tokens_ptr in place. The unconditional store to num_accepted_tokens_out_ptr at line 176 therefore compiles/writes through None on the v0.26.0 lane.
Please either backport the #50432 host-side snapshot/output-buffer flow to the v0.26.0 compatibility patch, or select the write contract explicitly at host/version specialization time. Using one unconditional output-pointer store for both lanes is not compatible.
- PRECOMPUTED_NEW_COMPUTED: restore running-state formula (new_num_computed - num_accepted + 1) instead of the aligned value. - Move num_scheduled/num_computed/num_draft loads into the else branch to avoid loading through None pointers on the precomputed path. - Validate batch_idx against num_reqs (not the mapped req slot) and handle the -1 skip sentinel under HAS_IDX_MAPPING. - Fix DS conv row copy: num_loops = dim_row_count, copy_size = (conv_width - bias) * elem_size, advancing src/dst by dim_row_stride. - Select the num_accepted write target by whether an output buffer is provided (main snapshot+out-buffer vs v0.26.0 in-place under V2). Signed-off-by: liaoqidan <1107297340@qq.com>
|
/rerun Failed:
|
1 similar comment
|
/rerun Failed:
|
…ernel dim_row_count/dim_row_stride were loaded inside the CONV_STATE_DIM_FIRST branch but referenced from the copy loop's CONV_STATE_DIM_FIRST and is_conv_state branch. Triton's control-flow scoping cannot prove the variable is defined across the narrower condition, so compilation failed with an undefined dim_row_stride. Load the DS row metadata inside the copy branch itself (same scope as the loop) and drop the num_loops intermediate, matching upstream's self-contained DS copy. Signed-off-by: liaoqidan <1107297340@qq.com>
|
/rerun Failed:
|
2 similar comments
|
/rerun Failed:
|
|
/rerun Failed:
|
6 participants
What this PR does / why we need it?
Upgrade baseline
2e09247c2d7b6b97d13af6e71a85bf8d1271deb6to58d3918e3ea0a544ffedadad2ba84559e9c51d8f. The full upstream range is available in this comparison.0.26.0compatibility lane while adapting the main lane to the new upstream contracts. Version gates usevllm_version_is("0.26.0")and are limited to real contract differences.Changes by file
1.
.github/vllm-main-verified.commit58d3918e2.
vllm_ascend/patch/platform/patch_fused_moe.pyFusedMoEtoFusedMoEFactory.FusedMoEFactory; on v0.26.0, also patch legacyFusedMoE.3.
vllm_ascend/models/deepseek_v4.py/vllm_ascend/models/minimax_m3/minimax_m3.py/vllm_ascend/ops/fused_moe/fused_moe.py/vllm_ascend/ops/fused_moe/routed_experts.pyFusedMoEFactory; remove deadFusedMoEre-exportFusedMoEwithFusedMoEFactory.FusedMoEre-export fromfused_moe.pyand dead reference inrouted_experts.pycomment.4.
tests/ut/models/test_deepseek_v4_moe.py/tests/ut/models/minimax_m3/test_minimax_m3.pyFusedMoEFactory"FusedMoE"→"FusedMoEFactory".5.
vllm_ascend/worker/model_runner_v1.pycalculate_kv_scalesremovalvllm_version_is("0.26.0")guard.clear_buffer()removalclear_buffer()fromRoutedExpertsCapturer.vllm_version_is("0.26.0")guard.6.
vllm_ascend/models/layer/attention/layer.pyQ/K/V_SCALE_CONSTANTreferencesq_range/k_range/v_rangeinitializations and deadimport envs.DSAAttention.forward()never used these attributes.7.
tests/ut/patch/platform/test_deepseek_v4_thinking.pylow/minimal/medium→low.vllm_version_is("0.26.0")guard.8.
tests/e2e/pull_request/one_card/spec_decode/test_extract_hidden_states.pychunked_prefillfor hybrid modelFalse→True.9.
vllm_ascend/patch/platform/patch_vision.py(new) +vllm_ascend/patch/platform/__init__.pyFusedInputNorm.forwardeps=0.0 → eps=1e-5FusedInputNormwithF.batch_norm(eps=0.0).eps=1e-5; guarded withcontextlib.suppress(ImportError).FusedInputNorm. Remove this patch once bundled PyTorch >= 2.13.0.10.
vllm_ascend/ops/triton/mamba/postprocess.pyCONV_STATE_DIM_FIRST,HAS_IDX_MAPPING,PRECOMPUTED_NEW_COMPUTED,state_dim_row_count/stride,idx_mapping_ptrparameters, andnum_loopsfor DS conv copy.11.
tests/e2e/conftest.py/tests/ut/spec_decode/test_speculators_vwn_eagle3.pymaybe_calc_kv_scalesmock12.
vllm_ascend/patch/__init__.pyCompatibility and review notes
vllm_version_is("0.26.0")exclusively; nohasattrfallbacks beyond the explicitly justifiedclear_bufferguard (where the upstream change is a method removal, not a rename).FusedMoE→FusedMoEFactoryrename is applied consistently across all call sites:deepseek_v4.py,minimax_m3.py,fused_moe.py,routed_experts.py, andpatch_fused_moe.py.layer.pyQ/K/V_SCALE_CONSTANTremoval is a dead-code cleanup: theDSAAttentionclass initialized these tensors fromenvsmodule-level constants (which still exist), but never used them inforward().Does this PR introduce any user-facing change?
No. This is a compatibility update; no new Ascend-specific public API is introduced.
How was this patch tested?
CI on the branch. See Buildkite workflow run for detailed results.