[CI] vllm main2main 0731 0351e9a - #13358
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 PR synchronizes vllm-ascend with the latest vLLM upstream changes while ensuring backward compatibility with vLLM 0.26.0. It introduces version-gated logic across core components, including distributed communicators, weight transfer engines, and custom ops, to handle upstream API evolutions. Additionally, it updates verified commit hashes and refactors the weight transfer engine to support stateful trainer-side operations. 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
|
There was a problem hiding this comment.
Code Review
Suggested PR Title:
[Ops][Misc] Add compatibility support for vLLM main branch (post-0.26.0)Suggested PR Summary:
### What this PR does / why we need it?
This pull request introduces comprehensive compatibility updates to ensure the `vllm-ascend` plugin works seamlessly with both vLLM `0.26.0` and the upstream `main` branch (post-`0.26.0`). It addresses various API, interface, and signature changes introduced in the upstream vLLM repository, including:
- Refactoring the NPU IPC weight transfer engine to support both the static `0.26.0` path and the stateful `IPCTrainerWeightTransferEngine` on `main`.
- Aligning signatures for `ParallelLMHead`, `NPUCommunicator`, `GroupCoordinator`, `MooncakeConnector`, `fused_moe` operators, and `InputBatch` to handle new arguments (such as `disable_tp`, `use_all2all`, `use_replayssm`, and `slot_mapping_modes`).
- Stubbing missing Triton modules (`triton.experimental.gluon`) and core functions (`_aggregate`) to prevent import failures on newer vLLM versions when using `triton-ascend 3.2.1`.
- Updating Triton kernels and Speculator/Eagle execution paths to match the updated upstream signatures (e.g., adding temperature/seeds parameters to `_prepare_dflash_inputs_kernel_ascend`).
Additionally, three critical issues were identified in the review:
1. In `examples/rl/rlhf_http_npu_ipc.py`, the local variable `client` shadows the global OpenAI `client`, which will cause an `AttributeError` during generation. It should be renamed to `sync_client`.
2. In `vllm_ascend/distributed/weight_transfer/npu_ipc_engine.py`, `device_index` is retrieved but not passed to `npu_generate_uuid()`, which can lead to generating the wrong physical NPU UUID.
3. The corresponding unit test in `tests/ut/distributed/weight_transfer/test_npu_ipc_engine.py` needs to be updated to assert that `npu_generate_uuid` is called with `device_index` on both branches.
### Does this PR introduce _any_ user-facing change?
No user-facing changes are introduced. These are internal compatibility and alignment updates to support different vLLM versions.
### How was this patch tested?
The changes are covered by existing and updated end-to-end and unit tests, which have been adapted to handle version-specific behaviors.|
👋 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. |
| if vllm_version_is("0.26.0"): | ||
|
|
||
| @triton.jit | ||
| def _prepare_dflash_inputs_kernel_ascend( |
There was a problem hiding this comment.
why change the ops? #48042 is not associate with it.
There was a problem hiding this comment.
The kernel change is required because upstream _prepare_dflash_inputs_kernel gained 4 temperature/seeds params in #50000 (Kimi K3); since the Ascend kernel is monkey-patched, the signature must match the upstream call site or Triton fails at launch. You're right that #48042 is unrelated — that was a typo, the comment/PR description will be corrected to #50000.
|
This pull request has conflicts, please resolve those before we can evaluate the pull request. |
4299ba0 to
8b4aee5
Compare
|
This pull request has conflicts, please resolve those before we can evaluate the pull request. |
Signed-off-by: hfadzxy <starmoon_zhang@163.com>
…nd tests Upstream vLLM PR #49665 changed VLLMValidationError to inherit from VLLMClientError instead of ValueError, breaking pytest.raises(ValueError) assertions that relied on the old inheritance relationship. Fix the 3 affected test cases to assert VLLMValidationError directly. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
- Add triton.experimental.gluon stubs to platform-level init before any
triton import, fixing ImportError(constexpr_type) in subprocesses such
as python -m vllm.model_executor.models.registry. Gated on
VLLM_VERSION != '0.26.0'.
- Add use_all2all kwarg to GroupCoordinatorPatch.__init__() gated on
!vllm_version_is('0.26.0') to match upstream GroupCoordinator signature
added in vllm main post-0.26.0. Ascend NPU has no all2all impl so the
parameter is only accepted for interface alignment.
Signed-off-by: hfadzxy <starmoon_zhang@163.com>
…ss coverage Placing gluon stubs in patch/platform/__init__.py only ran when pre_register_and_update() was called from EngineArgs.create_engine_config(), missing subprocesses such as python -m vllm.model_executor.models.registry where vllm.platforms discovery imports vllm_ascend but never calls pre_register_and_update(). Move stubs to module level in vllm_ascend/__init__.py so they execute the moment the plugin module is loaded - before any triton import chain can start. Also stub triton.language.core._aggregate, which vllm main also expects post-0.26.0 but triton-ascend 3.2.1 lacks. Revert stub placement from patch/platform/__init__.py. Revert conftest.py stub changes; keep a pointer comment to vllm_ascend/__init__.py. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
…InputBatch
Upstream InputBatch.__init__() added use_replayssm and slot_mapping_modes
parameters after 0.26.0. NPUInputBatch overrides __init__() without accepting
them, causing AttributeError on self.use_replayssm when the parent class
add_request() method runs. Accept the kwargs and store them under a
vllm_version_is('0.26.0') gate. NPU does not implement Mamba replay-SSM, so
the parameters are only accepted for interface alignment.
Signed-off-by: hfadzxy <starmoon_zhang@163.com>
- Add module-level triton gluon and _aggregate stubs to vllm_ascend/__init__.py
so they execute when the plugin module is imported by vllm.platforms
discovery, covering subprocesses such as python -m vllm.model_executor.
models.registry where pre_register_and_update() is never called. Gate
_aggregate import on find_spec('triton') to avoid ModuleNotFoundError on
310P and CPU-UT environments where triton is absent or mocked. Gated on
VLLM_VERSION != '0.26.0'.
- Add use_all2all kwarg to GroupCoordinatorPatch.__init__() gated on
!vllm_version_is('0.26.0') to match upstream GroupCoordinator signature
added post-0.26.0.
- Add use_replayssm and slot_mapping_modes kwargs to NPUInputBatch.__init__()
gated on !vllm_version_is('0.26.0') to match upstream InputBatch signature
added post-0.26.0. NPU does not implement Mamba replay-SSM; parameters are
only accepted for interface alignment.
- Revert conftest.py triton stubs; keep a pointer comment directing readers to
vllm_ascend/__init__.py.
Signed-off-by: hfadzxy <starmoon_zhang@163.com>
- Add disable_tp kwarg to AscendParallelLMHead and AscendParallelLMHead310
__init__() gated on vllm_version_is('0.26.0') to match upstream
ParallelLMHead signature added post-0.26.0. Ascend NPU uses its own TP
logic via lmhead_tp_enable(), so the parameter is accepted for interface
alignment only.
- Add used_output_is_reduced kwarg to AscendMoERunner.
_maybe_reduce_shared_expert_output() gated on vllm_version_is('0.26.0')
to match upstream signature. Ascend already handles shared expert
reduction in _forward_shared_experts, so the parameter is ignored.
Signed-off-by: hfadzxy <starmoon_zhang@163.com>
…tput
Upstream MoERunner._maybe_reduce_final_output() added output_is_reduced
parameter after 0.26.0. The call site at moe_runner.py:776 now passes 4
arguments (including fused_output_is_reduced), but AscendMoERunner's
override only accepted 3. Accept and ignore the kwarg under a
vllm_version_is('0.26.0') gate. Also widen trunc_size to int | None to
match the upstream signature.
Signed-off-by: hfadzxy <starmoon_zhang@163.com>
- test: expect VLLMValidationError instead of ValueError in structured output mixed backend tests (vllm#49665) - test: stub triton modules missing from triton-ascend 3.2.1 before vllm import in e2e conftest (vllm#48841) - patch: gate FLA imports with vllm_version_is(10.26.01) to avoid triton-ascend NPU driver crash during module load - spec_decode: gate prepare_inputs_to_capture keyword between skip_attn (v0.26.0) and full_cudagraph (main) (vllm#49364) - npu_communicator: accept use_all2all on main for interface alignment with upstream DeviceCommunicatorBase (vllm#47288) Signed-off-by: hfadzxy <starmoon_zhang@163.com>
Signed-off-by: hfadzxy <starmoon_zhang@163.com>
… vLLM version compat vLLM main (post-0.26.0) removed IPCTrainerSendWeightsArgs and the static trainer_send_weights path via: vllm-project/vllm#48042 — Stateful Trainer Send: New Abstractions [1/N] vllm-project/vllm#48981 — Stateful Trainer Send: IPC [2/N] Add a vllm_version_is conditional branch: - 0.26.0: preserves the existing static NPUIPCTrainerSendWeightsArgs + NPUIPCWeightTransferEngine.trainer_send_weights path unchanged. - main: introduces NPUIPCTrainerInitInfo (backend='npu_ipc'), NPUIPCTrainerWeightTransferEngine (subclass of upstream's IPCTrainerWeightTransferEngine), and delegates HTTP transport to HTTPVLLMWeightSyncClient. Register the trainer-side engine via WeightTransferTrainerFactory only on main. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
Signed-off-by: hfadzxy <starmoon_zhang@163.com>
…ter drift Upstream _prepare_dflash_inputs_kernel gained four new parameters (out_temperature_ptr, out_seeds_ptr, temperature_ptr, seeds_ptr) via vllm-project/vllm#50000. The Ascend kernel, monkey-patched in patch_v2/patch_triton.py, must present the same signature as the upstream call-site in prepare_dflash_inputs() to avoid 'got multiple values for argument' Triton errors at launch. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
vllm main added trainer_send_weights as an abstract method on WeightTransferEngine (30b4e7f479). Add a transitional stub to NPUIPCWeightTransferEngine (else branch only, not 0.26.0) to satisfy the ABC contract. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
…o for vllm main compat In vllm main, WeightTransferInitInfo no longer has packed field; it was moved to IPCWeightTransferInitInfo. Add packed: bool = False to NPUIPCWeightTransferInitInfo so init_transfer_engine can read it. Harmless in 0.26.0 where init_transfer_engine is a no-op. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
…ferEngine NPUIPCTrainerWeightTransferEngine.__init__ called super().__init__() which triggered torch.cuda.get_device_properties() in the parent IPCTrainerWeightTransferEngine.__init__, failing on NPU-only systems. Changes: - Bypass IPCTrainerWeightTransferEngine.__init__ to skip CUDA call; the NPU engine uses npu_uuid, not gpu_uuid, so this is safe. - Version-gate worker.update_weights: on 0.26.0 keep load_weights callback path; on main delegate to engine.update_weights() matching upstream gpu_worker.py pattern. - E2E test: register npu_ipc backend and close probe-initiated weight update so the stateful engine can drive its own lifecycle. - UT: fix dict access after asdict() serialization. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
Signed-off-by: hfadzxy <starmoon_zhang@163.com>
Signed-off-by: hfadzxy <starmoon_zhang@163.com>
8b4aee5 to
1be01b6
Compare
What this PR does / why we need it?
vllm_ascend/__init__.pyfrom triton.experimental import gluon,from triton.experimental.gluon import language as glandfrom triton.language.core import _aggregatetotriton_utils/__init__.py, requiring Triton 3.6+ API absent fromtriton-ascend 3.2.1triton.experimental.gluon/.languageassys.modulesstubs; stubtriton.language.core._aggregate. Gated onos.getenv("VLLM_VERSION") != "0.26.0", mirroringvllm_version_isenv var path.tests/ut/patch/platform/test_patch_structured_output.pyVLLMValidationErrorbase class fromValueErrortoVLLMClientError, breakingpytest.raises(ValueError)assertionspytest.raises(ValueError, ...)topytest.raises(VLLMValidationError, ...)tests/e2e/pull_request/one_card/test_guided_decoding.pyVLLMValidationErrorbase-class change; on 0.26.0 the upstream validation may still raiseValueErrorpytest.raises(ValueError, ...)on 0.26.0 vspytest.raises(VLLMValidationError, ...)on mainvllm_ascend/ops/vocab_parallel_embedding.pyvllm_ascend/_310p/ops/vocab_parallel_embedding.py*, disable_tp: bool = Falsekeyword arg toParallelLMHead.__init__()andVocabParallelEmbedding.__init__()__init__withvllm_version_is("0.26.0"): else branch acceptsdisable_tpkwarg and forwards it (Ascend manages TP vialmhead_tp_enable())vllm_ascend/distributed/kv_transfer/kv_p2p/mooncake_connector.pyvllm_ascend/distributed/kv_transfer/kv_p2p/mooncake_hybrid_connector.pyKVConnectorBase_V1.__init__()setsself._kv_transfer_config, required byrequires_kv_deliveryproperty__init__: else branch setsself._kv_transfer_configon vllm mainvllm_ascend/ops/fused_moe/fused_moe.pyfused_output_is_reducedkwarg to_maybe_reduce_shared_expert_output()andoutput_is_reducedkwarg to_maybe_reduce_final_output()trunc_sizemay now beNonevllm_ascend/patch/worker/patch_distributed.pyuse_all2all: bool = Falseparam toGroupCoordinator.__init__()_IS_VLLM_026env var check (avoids vllm-ascend import cycle). Else branch storesself.use_all2all(unused, no all2all on Ascend)vllm_ascend/distributed/device_communicators/npu_communicator.pyuse_all2all: bool = Falseparam toDeviceCommunicatorBase.__init__()NPUCommunicator.__init__withvllm_version_is("0.26.0"): else branch accepts and forwardsuse_all2all(NPU keeps the no-op_NpuAll2AllManager)vllm_ascend/worker/npu_input_batch.pyuse_replayssm: boolkwarg (Mamba replay-SSM); vllm#40996 — addedslot_mapping_modeskwarg (DCP hybrid attention)self.use_replayssmandself.slot_mapping_modesonly on main (not vllm_version_is("0.26.0")) for interface alignmentvllm_ascend/worker/v2/spec_decode/eagle/aclgraph.pyskip_attnparameter tofull_cudagraphin cudagraph mode logicskip_attn=(cg_mode != PIECEWISE); main passesfull_cudagraph=(cg_mode != PIECEWISE)vllm_ascend/distributed/weight_transfer/npu_ipc_engine.pyvllm_ascend/distributed/weight_transfer/__init__.pyIPCTrainerSendWeightsArgs/trainer_send_weightswith a statefulIPCTrainerWeightTransferEnginedriven byWeightTransferTrainerFactory.trainer_init(...).send_weights();packedmoved onto trainer init-info; per-roundupdate_infoslimmedNPUIPCWeightTransferEngine; main defines statefulNPUIPCTrainerWeightTransferEngine(_send_unpackedinstance method,packedonNPUIPCWeightTransferInitInfo, no-argnpu_generate_uuid(),is_sender/_all_gather_and_merge_handles).register_engine()also registers the trainer engine inWeightTransferTrainerFactoryon main onlyexamples/rl/rlhf_http_npu_ipc.pyWeightTransferTrainerFactory.trainer_init(...)+engine.send_weights()withHTTPVLLMWeightSyncClient/ModuleSourceNPUIPCWeightTransferEngine.trainer_send_weights; main uses the stateful trainer engine pathtests/e2e/pull_request/one_card/test_npu_ipc_weight_transfer.py_post(start/finish)+ statictrainer_send_weights; main drivesWeightTransferTrainerFactory.trainer_init(...).send_weights()(engine owns the lifecycle)tests/ut/distributed/weight_transfer/test_npu_ipc_engine.pyIS_VLLM_026):_send_unpacked,packedplacement,npu_generate_uuidarity, andstart/finish_weight_updateno-op behavior asserted per versionvllm_ascend/worker/v2/spec_decode/dflash/speculator.py_prepare_dflash_inputs_kernelgainedtemperature/seedssampling params for probabilistic draft sampling_prepare_dflash_inputs_kernel_ascend: 0.26.0 branch keeps the old signature; main branch adds the fourtemperature/seedspointers and the corresponding stores to stay aligned with upstream while keeping Ascend's own kernel.github/vllm-main-verified.commit0351e9aa1fdf1a51329d1906881528dfe61fc88eDoes this PR introduce any user-facing change?
How was this patch tested?