Repository navigation
build: move the MUSA stack to torch 2.11.0.post2, torchada 0.1.90 and deep_ep post2 - #240
Conversation
3826cea to
8b724d9
Compare
yeahdongcn
left a comment
There was a problem hiding this comment.
Merge verdict: no blocking findings — ready to merge.
Reviewed head 8b724d97d1b7bb7747cf0beed9460458ac6b7f51 (identical to the dispatch hint, so no force-push between dispatch and review) against base v0.28.0-dev @ 7717608009; draft: true, mergeable_state: clean, 2 commits, 9 files, +27/-24. gh pr checks 240 --repo MooreThreads/vllm-musa reports no checks reported on the 'codex/nemotron35-tpot-20260918' branch, and the head tree contains no .github/workflows, so nothing automated ran: everything below is static/metadata verification plus the PR's own on-card probes. Approval is deferred to the terminal round per this round's policy.
Findings
P2 — stale rationale and an internal tracker ID in shipped source — vllm_musa/model_executor/layers/fused_moe/router/musa_router_gemm.py:70-73
* ``torch.mm(..., out_dtype=torch.float32)`` on this MUSA stack returns an all-zero
tensor (see MUSA-100042): the GEMM never writes its output, for BF16 *and* FP32
inputs. Any FP32-router path that relies on ``out_dtype`` on MUSA is silently
wrong, which is a further reason to compute the gate in this kernel instead.This PR is what invalidates the claim. The branch now pins torch_musa 2.11.0.post2 + torchada 0.1.89, and 0.1.89 arms its torch.mm/torch.bmm out_dtype backport for every torch_musa < 2.13.0 (torchada/_patch.py:576-597: if version_of(...) >= "2.13.0": return, otherwise it wraps both ops and calls _register_jit_builtin_alias). The PR's own probe (c) records that exact call going from all zeros, abs_sum 0.0 to max_abs_err 9.54e-07, abs_sum 769.98 == ref. So on the pinned stack an out_dtype= call is no longer silently wrong, and "a further reason to compute the gate in this kernel instead" no longer follows; the kernel's remaining justification is the perf table. The same block cites tracker MUSA-100042, which comment-style.md rule 2 forbids in shipped comments/docstrings/commit messages, and line 25 records the measurement environment as the old vllm:v0.28.0-ph1-5.2.0-torch2.11.0.post1-latest image. The same ticket ID is also in the message of commit bdd577ab9905.
Suggested: rewrite those two sentences to state what holds on the pinned stack (repaired below torch_musa 2.13.0), keep the measurement as the module's reason for existing, re-label or re-measure the provenance line against the post2 image, and drop the MUSA-100042 reference from both the docstring and the commit message.
P2 — no automated or end-to-end run covers the stack switch — requirements/common.txt:2, requirements/musa_private.txt:15-18
Trigger: any serving process on the pinned stack. This change is not only version strings — four torchada shims switch off at post2, so torch.multinomial, Tensor.log_, torch.accelerator.memory_* and inductor MUSA template registration all change implementation together (torchada/_cpp_ops.py:105-107 sets TORCHADA_DISABLE_OP_OVERRIDE_{multinomial,log,log_}=1; torchada/_patch.py:145, :191, :2294 gate the rest). vllm_musa/v1/sample/topk_topp_sampler.py:194 calls torch.multinomial directly behind a gate (can_use_musa_seeded_multinomial, :103-116) that was validated on the previous stack, and the PR's multinomial probe only checks the device of the returned picks.
Not blocking by itself, because the PR body scopes the model matrix to the tracked follow-up MUSA-100018 and the pin metadata checks out (below). Before the pins become the shipped default, or as the first item of that follow-up: run the in-tree slices that exercise the switched ops on the post2 image (pytest tests/test_chunked_min_p_sampler.py tests/test_chunked_min_p_sampler_source.py tests/test_gdn_fused_zba_compile.py tests/test_fused_add_rms_norm_compile.py -q) plus one serving smoke with a real request. I did not run them: they need torch/MUSA. Related check, not a demonstrated defect: MUSA-runtime torch.mm/torch.bmm calls that pass out_dtype= now hit torchada's wrapper, whose deferred vendor probe fires on the first such call; the repo's own DeepSeek-V4 helper avoids that path (vllm_musa/patches/series/0014-MUSA-vllm.models.deepseek_v4.attention.patch uses F.linear(a.to(out_dtype), weight.to(out_dtype)) on MUSA), so I could not confirm any live caller.
P0: none. P1: none.
Verified (read at the head commit)
- Provider metadata, re-read live:
torch,torch_musa,torchvisionandtorchaudioall publish…post2+musa5.2.0wheels forcp310andcp312linux_x86_64inhttps://dl.mthreads.com/repo/api/pypi/pypi/simple/<pkg>/; the only newer-looking artifacts are the separate aarch64+musa5.2.1mp22line, still at post1, as the PR body states. deep_ep: the index has notorch2.11.0.post2build. I downloaded the pinned cp310 wheel (74,831,967 bytes) and read itsMETADATA:Name/Versiononly, noRequires-Distat all, so Docker step 3 cannot resolve it to a different torch. The new comment inrequirements/musa_private.txtis accurate.torchada==0.1.89is the newest PyPI release and declares onlyRequires-Dist: torch. From its published sdist,torchada/utils/cpp_extension.pystill definesstable_compat_box_header,stable_compat_include_dirandinclude_paths, whichsetup.py:312-320imports at build time — the 0.1.89 internal refactor does not break this repo's build path.- The torchada gates match the PR body: four shims installed only below post2 (
_patch.py:132-146,:191,:2294,_cpp_ops.py:105) and themm/bmmbackport armed below 2.13.0. - Whole-head-tree scan (517 files, streamed from the head tarball) for
2.11.0.post1,torchada==0.1.8[0-8]andMATE 0.2.6: only intentional leftovers (thedeep_epdistribution name, the untouched 2.9.1 stack entry, the pre-0.2.6 MATE patch labels) plus the docstring in the first finding. README/README_CN,docker/README.mdand both installation docs agree with the pins, and the documented branch pinsmate==0.2.7/tilelang_musa==0.1.12matchrequirements/musa_private.txt. docker/musa.Dockerfile:350-353: the reorder satisfiestests/test_inplace_porting_contract.py:123-126. I reconstructed the embedded script from the head Dockerfile and compiled it (SYNTAX: OK); the firstimportlib.import_module(module_name)precedes the singleinstalled = version(dist_name),installedis still assigned on every iteration so the threeskip_importdistributions keep their version check, and the embedded script has not moved. In-memory re-run of that test's pin assertions against the head file contents (stack selection, private subset, torchada-in-common.txt, exact-gate strings, both ordering assertions) all hold.- Image tags:
docker/build_image.sh:22-30derives the tag fromrequirements/musa_private.txt, so2.11.0.post2+musa5.2.0becomes2.11.0.post2_musa5.2.0— matching both updated occurrences indocker/README.md. - The new 3-line comment cannot leak into
install_requires:setup.py:58andtests/test_mate_027_dependency_contract.py:31both skip#lines. The claimed19 passedis consistent with 17 + 2 test functions in the two named files.
Not verified / assumed
- No image build,
pytestrun, GPU probe or serving smoke was executed in this session; every runtime number quoted above is the PR author's, reproduced nowhere by me. - The published-image claim in
docs/installation.md/docs/installation-cn.md(image reports MATE/Flash 0.2.4, tilelang_musa 0.1.8) is untouched by this PR and was not re-checked against the registry. pip checkcleanliness fordeep_epis inferred from the wheel carrying no requirements at all, not from running pip against the new stack.- One nuance in the PR body: "0.1.89 is the release that turns off four of its own legacy shims at post2" overstates the history — 0.1.83 already gated the
torch.acceleratorshim on_TORCH_MUSA_ACCELERATOR_FIX_VERSION = "2.11.0.post2". What 0.1.89 adds for this pairing is themultinomial/log/log_gating and themm/bmmbackport. No code impact, but worth not repeating the stronger claim in shipped docs.
auto-review round 1/3 · verdict: clean
8b724d9 to
d1f588b
Compare
Move the v0.28.0-dev MUSA stack to the newest published release train (verified against the Moore Threads index on 2026-09-24): torch 2.11.0.post1+musa5.2.0 -> 2.11.0.post2+musa5.2.0 torch_musa 2.11.0.post1+musa5.2.0 -> 2.11.0.post2+musa5.2.0 torchvision 0.26.0.post1+musa5.2.0 -> 0.26.0.post2+musa5.2.0 torchaudio 2.11.0+musa5.2.0 -> 2.11.0.post2+musa5.2.0 torchada 0.1.83 -> 0.1.89 torchvision and torchaudio move with torch because the vendor published all four wheels together on 2026-09-23 as one post2 train. deep_ep stays on its newest published build: no 2.11.0.post2 wheel exists and the wheel declares no torch requirement, so the pin is annotated instead of silently lagging. torchada 0.1.89 is the release that gates its legacy shims on torch_musa >= 2.11.0.post2, so it is the intended partner of this torch bump; it also carries the torch.mm/bmm out_dtype backport for the zero / wrong-value result reported in MUSA-100042. Also corrects the documented MATE version (0.2.6 -> 0.2.7), which had been stale since the MATE 0.2.7 cohort landed.
0.1.90 touches one file. _discover_factory_functions() now also returns the subset of factories torch itself device-injects into, and only that subset is added back to torch.utils._device._device_constructors(). The *_like family is no longer registered there - that registration was what made torch.zeros_like(x) return a cpu tensor whenever a device context was active (torch.set_default_device pushes one even for cpu), feeding host pointers into device kernels. It also re-lays out the fused-MoE tuning files: the 41 duplicated 'block_shape=[128,128]' (space-free) names are gone and the runtime now aliases them onto the canonical 'block_shape=[128, 128]' files, which is why the wheel is smaller. Declared dependencies, Requires-Python and the E=128 shape coverage are byte-for-byte unchanged, and the E=128,N=1856 Nemotron config is identical.
deep_ep was the last component in the post2 stack still built against post1, and the only one with a real ABI question: its wheel declares no torch requirement at all, so it installs beside the post2 torch without the resolver noticing the skew. 1.1.0+musa5.2.0torch2.11.0.post2 is now on the index (wheel sha256 6a7973345617..., matching the index's own declaration). Measured against the post1 wheel, the only file that changed is deep_ep_cpp.cpython-310-x86_64-linux-gnu.so; every Python source is byte-identical. That extension does change shape: it adds libmusa_kernels.so, libdl, librt and libmlx5, and its RUNPATH now points at the builder's /opt/_internal prefix instead of this image's dist-packages. Neither is a problem here - libmusa_python.so.2 already depends on libmusa_kernels.so, so it is in the process for any torch_musa import, and ldd resolves it from torch_musa's own lib dir. Verified by installing the wheel into the post2 image and importing deep_ep and deep_ep.utils.
yeahdongcn
left a comment
There was a problem hiding this comment.
Merge verdict: no blocking findings — ready to merge. Verified at head d1f588bfdae3: no earlier P0/P1, no CI configured, and this round's delta (torchada 0.1.90, deep_ep post2) holds against the live index and the wheels themselves. A clean verdict ends auto-review here — a human can merge.
Findings: 0 blocking · 3 nit
Earlier: round 1 (8b724d97) raised no P0/P1; its two P2 items are rows 1–2, still present at this head.
| nit | location | problem | fix |
|---|---|---|---|
| nit | vllm_musa/model_executor/layers/fused_moe/router/musa_router_gemm.py:71 |
out_dtype "all-zero on this stack" claim is now false; the tracker ID is forbidden in shipped source |
say "repaired below torch_musa 2.13.0"; drop MUSA-100042 |
| nit | …/router/musa_router_gemm.py:25 |
provenance names the superseded …torch2.11.0.post1-latest image |
re-measure, or re-label, on the post2 image |
| nit | requirements/common.txt:2 |
41 of the 100 space-free fused-MoE config names are symlinks created at import torchada; a read-only site-packages drops them silently |
assert the config-name count in the image's build gate |
What changed · what I could not verify
Diff map (8 files, +25/−25): the only runtime-bearing lines are the six pins in requirements/musa_private.txt and requirements/common.txt; tests/test_inplace_porting_contract.py moves the stack table; README/README_CN, docker/README.md and both installation docs are prose. Nothing else in the tree changed.
Round-1 → this head. torchada 0.1.89 → 0.1.90, deep_ep post1 → post2, MATE prose 0.2.6 → 0.2.7. The post2 move of the four torch packages was already at round 1's head; round 1's deep_ep: the index has no torch2.11.0.post2 build is now resolved, and its torchada==0.1.89 objection is moot because the pin moved to 0.1.90.
Verified live, this session
- The MT index carries
post2+musa5.2.0x86_64 wheels forcp310andcp312fortorch,torch_musa,torchvisionandtorchaudio; the only post1 leftovers are the separateaarch64 +musa5.2.1mp22line. deep_eppost2 cp310 wheel: 74,953,528 bytes, sha2566a7973345617d847…— the index fragment and the PR body agree. ItsMETADATAdeclaresName/Versiononly, noRequires-Distat all, so no resolver guards this pin; onlydeep_ep_cpp.cpython-310-x86_64-linux-gnu.sodiffers from post1 (every other entry is CRC-identical, so the "Python sources byte-identical" claim holds). That.so'sDT_NEEDEDislibmusa_python.so.2, libmusa_kernels.so, libtorch_python.so, libmusart.so.5, libdl.so.2, librt.so.1, libmlx5.so.1, libstdc++.so.6, libm.so.6, libgcc_s.so.1, libc.so.6, ld-linux-x86-64.so.2. The four the body singles out are present; that they resolve inside the image is the author's evidence, not mine — I cannot read the image.torchada0.1.90 is the newest PyPI release and declares onlyRequires-Dist: torch.torchada/utils/cpp_extension.pystill definesstable_compat_box_header,stable_compat_include_dirandinclude_paths, whichsetup.py:312-320imports at build time, so 0.1.90 does not break this repo's extension path.- The 0.1.89 → 0.1.90 source delta is exactly what the body claims:
_patch.py's_discover_factory_functions()now returns(names, device_injectable)and the_device_constructors()re-add is restricted todevice_injectable, so the*_likefamily keeps input-derived device; plus the fused-MoE config re-layout, version bumps,csrc/ops.h, and new regression tests. The other gates still target this pairing:_cpp_ops.py:105disables themultinomial/log/log_overrides attorch_musa >= 2.11.0.post2,_patch.py:582arms themm/bmmout_dtypebackport below 2.13.0. - The config re-layout loses nothing. 0.1.90 ships 100 config files (59 space-free, 41 spaced);
_alias_space_free_names()runs atimport torchada(called fromtriton/autotune/fused_moe/__init__.py, just before it setsVLLM_TUNED_CONFIG_FOLDER), giving 100 space-stripped lookup names — identical to 0.1.89's 100 and a superset of 0.1.83's 98. No tuned shape changes target, which is the one thing a green MoE smoke could not have shown. - Contract tests re-run in memory against the head file contents: 19/19 assertions hold (single
torch==pin, supported stack key, private subset,torchada==0.1.90incommon.txt, the six-distexact_version_distsliteral, torchada-gate-before-torch-gate and import-before-version ordering,triton==3.2.0). The head'smusa_private.txtis 16 unique exact pins with no duplicates, sotest_mate_027_dependency_contract._exact_requirements()parses it. - A whole-tree scan (525 files) finds no
2.11.0.post1or0.1.83-era reference other than rows 1–2; docs agree with the pins (build_image.shderives the tag from the pin, giving…_2.11.0.post2_musa5.2.0).
Not verified / assumed
- No image build, no
pytest, no GPU run, no CI. The tree has no.github/directory, andgh pr checks 240 --repo MooreThreads/vllm-musareturnsno checks reported on the 'codex/nemotron35-tpot-20260918' branch(exit 1), so nothing automated can run on this branch. Every runtime number in the PR body is the author's and reproduced nowhere by me. - Unconfirmed, would need the post2 image:
vllm_musa/v1/sample/topk_topp_sampler.pycallstorch.multinomialbehindcan_use_musa_seeded_multinomial(), and torchada's override for it switches off at post2 in favour of the native op. The gate is not env-driven, so I cannot tell statically whether the legacy per-request-seed path still behaves; a seeded sampling smoke on the post2 image would confirm it. - The body's row-4 MLA failure (
'ChunkedContextMetadata' object has no attribute 'seq_tot') is fork code absent from this diff; I take its MUSA-100055 attribution on trust rather than having reproduced it.
auto-review round 2/3 · verdict: clean
d1f588b to
078be0e
Compare
…gate) At `v0.28.0-dev` the documented regeneration path was dead. `musa_sync.py rebase` (== `make -f Makefile.sync apply-patches`) stops at `0014-...` with `corrupt patch at line 78`, and after repairing that entry it stops at `0140` with `Patch format detection failed`. The build and `verify` stayed green because `build_apply.py:51-60` applies entries with `git apply --recount -p1`, which infers the hunk counts from the patch body and accepts bare diffs, so only the `git am` path noticed. * three entries carried hunks whose header disagrees with their body: `0014`'s `@@ -6,58 @@` hunk carries 60 lines (and the later offsets are off by two as a result), plus one hunk each in `0133` and `0138`. `0014` is restored from the commit it was generated from; every hunk in the series now matches its header. * twelve entries were bare diffs rather than mailboxes (`0140`-`0150`, `0166`), which `git am` refuses. Every entry is a mailbox again, renumbered contiguously. The rewritten series is a fixed point: at the pre-rebase 171-entry head it applied to the pinned upstream (`2cf0a6915ce544dc493a0990f2ea38d81601128a`) to `f09964462be438ffdb982ebf018b1692ac55448c` (171 entries, 138 changed files), the identical tree the pre-rewrite series produced, and `git am -3` reproduced it. After the rebase onto `8acf64368b` the same fixed point holds for the 172-entry series: the base's series and this one, each applied to a pristine pin, give the identical sha1 tree `abc996438a3ba0291dbf19429e60a3924f5f6740` - verified with the build applier, with plain `git apply --recount -p1` + `git write-tree`, and with a `git am` replay of all 172 entries that lands on the same tree, the two applied working directories being identical file for file (ignored files included). The change is therefore form only; no patched line of upstream vLLM moves. `tools/musa_sync.py` gains `check-series` and a hardened `verify`/`--replay`/`--round-trip`: mailbox format, contiguous numbering, the manifest census, every gitlink shape, and a verdict line on every path that has one. `tools/patch_validate.py`, `Makefile.sync`, the developer guide and the tests follow. Measured on this tree: `check-series` PASS (172 clean / 172 total), `--replay` and `--round-trip` PASS in the states they document, `verify` `174 clean / 210 total / 1 need attention` - that single row is the base's own pre-existing `module-drift` tripwire for `v1.attention.backends.mla.common`, and `regen --area module` clears it (175 clean / 210 total) - and `pytest tests/test_musa_sync.py` -> 104 passed (122 across the gate's three test files). A from-scratch image built from this tree is byte-identical to it on all 56 changed paths, and the 8-leg model smoke matrix on that image repeats `pass=5 fail=2 skip=1`, leg for leg as before; both are recorded in the PR description. Rebased onto 8acf643 (#240/#260/#263) and re-derived: the series is 172 entries (the base's 0172 is re-emitted by `regen`, so it is no longer inherited verbatim). `tests/test_musa_sync.py` no longer pins the entry count - it derives it from the series directory - because a fixed count reddens the gate on every base-side patch addition, which is exactly the coupling this PR removes. Re-measured on the 172-entry series, the declared-blob census the gate's docstrings cite is 232 postimages / 232 preimage pairs (230 at 171 entries): 0 postimages and 133 preimages resolve in a pin-only checkout, and 232/232 postimages are the declared path's blob in the entry's own commit - so the docstrings now state the measured numbers.
…gate) (#250) At `v0.28.0-dev` the documented regeneration path was dead. `musa_sync.py rebase` (== `make -f Makefile.sync apply-patches`) stops at `0014-...` with `corrupt patch at line 78`, and after repairing that entry it stops at `0140` with `Patch format detection failed`. The build and `verify` stayed green because `build_apply.py:51-60` applies entries with `git apply --recount -p1`, which infers the hunk counts from the patch body and accepts bare diffs, so only the `git am` path noticed. * three entries carried hunks whose header disagrees with their body: `0014`'s `@@ -6,58 @@` hunk carries 60 lines (and the later offsets are off by two as a result), plus one hunk each in `0133` and `0138`. `0014` is restored from the commit it was generated from; every hunk in the series now matches its header. * twelve entries were bare diffs rather than mailboxes (`0140`-`0150`, `0166`), which `git am` refuses. Every entry is a mailbox again, renumbered contiguously. The rewritten series is a fixed point: at the pre-rebase 171-entry head it applied to the pinned upstream (`2cf0a6915ce544dc493a0990f2ea38d81601128a`) to `f09964462be438ffdb982ebf018b1692ac55448c` (171 entries, 138 changed files), the identical tree the pre-rewrite series produced, and `git am -3` reproduced it. After the rebase onto `8acf64368b` the same fixed point holds for the 172-entry series: the base's series and this one, each applied to a pristine pin, give the identical sha1 tree `abc996438a3ba0291dbf19429e60a3924f5f6740` - verified with the build applier, with plain `git apply --recount -p1` + `git write-tree`, and with a `git am` replay of all 172 entries that lands on the same tree, the two applied working directories being identical file for file (ignored files included). The change is therefore form only; no patched line of upstream vLLM moves. `tools/musa_sync.py` gains `check-series` and a hardened `verify`/`--replay`/`--round-trip`: mailbox format, contiguous numbering, the manifest census, every gitlink shape, and a verdict line on every path that has one. `tools/patch_validate.py`, `Makefile.sync`, the developer guide and the tests follow. Measured on this tree: `check-series` PASS (172 clean / 172 total), `--replay` and `--round-trip` PASS in the states they document, `verify` `174 clean / 210 total / 1 need attention` - that single row is the base's own pre-existing `module-drift` tripwire for `v1.attention.backends.mla.common`, and `regen --area module` clears it (175 clean / 210 total) - and `pytest tests/test_musa_sync.py` -> 104 passed (122 across the gate's three test files). A from-scratch image built from this tree is byte-identical to it on all 56 changed paths, and the 8-leg model smoke matrix on that image repeats `pass=5 fail=2 skip=1`, leg for leg as before; both are recorded in the PR description. Rebased onto 8acf643 (#240/#260/#263) and re-derived: the series is 172 entries (the base's 0172 is re-emitted by `regen`, so it is no longer inherited verbatim). `tests/test_musa_sync.py` no longer pins the entry count - it derives it from the series directory - because a fixed count reddens the gate on every base-side patch addition, which is exactly the coupling this PR removes. Re-measured on the 172-entry series, the declared-blob census the gate's docstrings cite is 232 postimages / 232 preimage pairs (230 at 171 entries): 0 postimages and 133 preimages resolve in a pin-only checkout, and 232/232 postimages are the declared path's blob in the entry's own commit - so the docstrings now state the measured numbers. Co-authored-by: Xiaodong Ye <yexiaodong60@mthreads.com>
Rebased onto
v0.28.0-dev@0fae3a1948— the base now carries the MLA chunked-context migration(#260) — and moved to the newest published MUSA release train:
torch/torch_musa2.11.0.post2+musa5.2.0,torchada0.1.90, anddeep_ep1.1.0+musa5.2.0torch2.11.0.post2.The previous head (
8b724d97d1) also carried a Dockerfile import-order fix. The new base alreadycontains it — upstream #262 reordered the same dependency-gate loop and extended it to
transformers/huggingface_hub— so the rebase dropped that commit as already-upstream and thisPR is now purely the pin move plus its contract-test and documentation updates.
torch2.11.0.post1+musa5.2.02.11.0.post2+musa5.2.0torch_musa2.11.0.post1+musa5.2.02.11.0.post2+musa5.2.0torchvision0.26.0.post1+musa5.2.00.26.0.post2+musa5.2.0torchaudio2.11.0+musa5.2.02.11.0.post2+musa5.2.0torchada0.1.830.1.90*_likedevice-injection fix (below)deep_ep1.1.0+musa5.2.0torch2.11.0.post11.1.0+musa5.2.0torch2.11.0.post2mate,mate-mubin,flash_attn_3,flash_mla,deep-gemm,flashinfer-python,sageattention0.2.7tilelang_musa,triton,apache-tvm-ffi,torch-c-dlpack-extPrevious head
8b724d97d1→ this revision078be0e7c2.With this change the image no longer contains any component left on the post1 stack.
Is this the newest?
Queried the Moore Threads index directly on 2026-09-29 rather than trusting memory or an older
ticket:
cp310,linux_x86_64)torch2.11.0.post2+musa5.2.0torch_musa2.11.0.post2+musa5.2.0torchvision0.26.0.post2+musa5.2.0torchaudio2.11.0.post2+musa5.2.0torchada0.1.90deep_ep1.1.0+musa5.2.0torch2.11.0.post2torchvision/torchaudiomove withtorchrather than trailing a revision.post2is also built forcp312, so it is not a Python-version-specific artifact.aarch64+musa5.2.1mp22builds, a separate platform line still atpost1.deep_eppost2 — the pin that could not move in the previous revision — is now published:deep_ep-1.1.0+musa5.2.0torch2.11.0.post2-cp310-cp310-linux_x86_64.whl,sha256
6a7973345617d847…, 74 953 528 bytes (the index's own declared size matches the file).Why the
torchadabump belongs with thistorchbumptorchada0.1.83→0.1.90is a small series but0.1.90is the release that fixes asilent correctness bug in device injection:
_discover_factory_functions()used to re-add every device constructor as anout=-style factory, including the*_likefamily. With any device context active(
with torch.device("cpu"),torch.set_default_device("cpu")),torch.zeros_like(x)on a MUSAtensor therefore returned a CPU tensor. The downstream kernels took that host pointer as a
device address — silently, with no error at the call site; the upstream comment names sgl-kernel
dereferencing it.
0.1.90restricts the re-add to the original_device_constructors()set(
device_injectable), so the*_likefamily keeps its input-derived device while explicitdevice=calls are still honoured.This is why the
torchadapin moves together with thetorchbump instead of being left at0.1.83:0.1.89also removes four legacy shims thattorch_musa >= 2.11.0.post2implementsnatively (inductor MUSA template heuristics, float64 in-place
Tensor.log_,torch.acceleratormemory overrides, compiled
multinomial/log/log_) and arms thetorch.mm/torch.bmmout_dtypebackport belowtorch_musa 2.13.0(the runtime fix for MUSA-100042).0.1.90keepsall of that.
0.1.90also re-lays-out its bundled Triton fused-MoE configs:configs/triton_3_2_0/ships 100files instead of 141, because 41 space-free duplicate names are now created as aliases at import
time (
_alias_space_free_names()), which is what vLLM'sget_config_file_name()needs because itstrips spaces. Behaviour-preserving, but a green MoE smoke does not prove it — a missing config
silently degrades to defaults — so this PR verifies resolution directly (probe B below).
Why
deep_epmoves nowThe post2 wheel exists as of 2026-09-29. Against post1 the Python sources are byte-identical and
only
deep_ep_cpp.cpython-310-x86_64-linux-gnu.sodiffers: it gainsNEEDEDlibmusa_kernels.so,libdl.so.2,librt.so.1,libmlx5.so.1and dropslibpython3.10.so.1.0.Neither wheel declares any
Requires-Dist, so no resolver would have caught an ABI mismatch —hence it was checked by hand rather than assumed:
libmusa_kernels.soundertorch_musa/lib/,libmlx5.so.1in/usr/lib/x86_64-linux-gnu/,libmusart.so.5in/usr/local/musa-5.2/lib/), andlddresolves them;libmusa_python.so.2already depends onlibmusa_kernels.so, so it is mapped by anytorch_musaimport;deep_ep1.1.0+musa5.2.0torch2.11.0.post2,Buffer present: True.The build's own dependency gate now imports
deep_eppost2 as part of the image build, so afuture break here fails the build instead of a model run.
Validation
Static
pytest tests/test_inplace_porting_contract.py tests/test_mate_027_dependency_contract.py -q→ 20 passed (the contract test asserts the post2 stack key and the
torchada==0.1.90pin).tests/importstorch/numpyat module scope and is not collectible in the reviewinterpreter (
ModuleNotFoundError); this diff touches no runtime Python path, so those modulesare
not runlocally rather than reported as green.Image rebuild — the canonical path, from this branch
pr240-rebase@078be0e7c285b1171d424f7f9ce1d757f78d4f6f, clean tree.dockerignoreexcludesthird_party/vllmandsetup.pyclones the pinned upstream revision inside the image —which is also why this rebuild cannot be done from a cluster node
13:29:19→13:42:49(+08:00), 13m29s. Shorter than a cold build because thedependency layers were already in the build cache; the rebase only invalidated the application
layers.
requirements/musa_private.txt787cc560…,requirements/common.txt450c6fbf…,third_party/PINSc2ac34a4…(VLLM_TAG=v0.28.0,VLLM_COMMIT=2cf0a6915ce544dc493a0990f2ea38d81601128a)vllm-musa:ubuntu22.04_py3.10_musa_runtime_5.2_pytorch_release_2.11.0.post2_musa5.2.0,id
sha256:b93ba03882cb…, 20.8 GB. The tag is derived fromrequirements/musa_private.txt, sothe tag itself is evidence the pin took effect.
registry.mthreads.com/mcconline/inference/vllm:v0.28.0-ph1-5.2.0-torch2.11.0.post2-20260929-r2,digest
sha256:ee6c919aee57…; the node re-pull yields the identical image id.5b8d95a4a9, i.e. before the base gained fix(mla): migrate MUSA MLA chunked-context prefill to the vLLM v0.28 chunk layout #260) wasbuilt and published as
…-20260929, image idsha256:7ef2ae44fc5e…, digestsha256:478c89311b54….It is kept for traceability because its matrix is the one that exposed the now-fixed MLA row.
org.opencontainers.image.revision=078be0e7c285b1171d424f7f9ce1d757f78d4f6f,com.mthreads.vllm.version=v0.28.0The dependency gate is the image's own build step, so this is build-log output:
and the pin resolution itself:
Model matrix and targeted probes — S5000, host
10.20.36.6, GPU 0, driver5.2.0-serverA container from the image pushed above,
MUSA_VISIBLE_DEVICES=0, one card. Preflight recorded inthe same summary as the rows: one visible
MTT S5000, uuidd6a04b9f-fec2-2c87-aa03-a8360298e63b,torch/torch_musa2.11.0.post2+musa5.2.0,torchada0.1.90. The container's visible deviceuuid equals the lease's uuid, so the card scoping was verified rather than assumed.
Every row is a functional smoke: rows 1-4 run
--enforce-eagerand each row asserts theattention backend from vLLM's own
Using <X> backend.log line rather than from the environmentvariable we set. The Nemotron row uses the frozen compiled protocol from MUSA-100035 but sends a
single request and records no timing. No TPOT, acceptance-length or throughput claim is made
anywhere in this PR.
Qwen3-8B(bf16)FLASH_ATTNPASS … contains_Beijing=True backend_asserted=FLASH_ATTNQwen3-30B-A3B-FP8FLASH_ATTN+ fused MoE JITPASS qwen3_moe_fp8 quantization=fp8 contains_Beijing=True backend_asserted=FLASH_ATTNDeepSeek-V2-Lite-ChatFLASHMLAPASS dsv2lite_mla contains_Beijing=True backend_asserted=FLASHMLAvllm serve, 5-turn chatFLASHMLAPASS dsv2lite_server_multiturn requests=5 multi_turn=True backend_asserted=FLASHMLANVIDIA-Nemotron-3.5-Lightning-30B-A3B-BF16FLASH_ATTN+ MTPPASS nemotron35_mtp6 served=True mtp=True num_speculative_tokens=6 canonical_protocol=True eager=False functional_smoke_only=TrueTwo targeted probes cover what the matrix cannot (both
PASS):*_likedevice injectionzeros_like/ones_like/rand_likeon a MUSA tensor stay on MUSA under a CPU device context, without breaking explicitdevice="cpu"?ctx_cpu_zero_like_is_musa=True,default_cpu_*_like_is_musa=True(all three),explicit_cpu_still_cpu=Truelookups_pass=True;VLLM_TUNED_CONFIG_FOLDER=…/torchada/triton/autotune/fused_moe/configs/triton_3_2_0; 100 space-free + 41 spaced = 141 files, aliases byte-identical to their targetsThe Nemotron row loads the 61.4 GiB BF16 checkpoint, reports
Detected MTP model. Sharing target model embedding weights with the draft model., compiles itsMUSA MoE and norm JIT kernels, warms up the Mamba2 SSD Triton kernels, captures 13 decode graphs
and answers a chat completion.
Its first run used a 24-token budget, which the model spent entirely on a reasoning preamble
(
Here's a thinking process: 1. **Analyze User Input:** …), so the substring check onBeijingread
Falseeven though serving was correct. Re-run on this same image with the identical protocoland a 64-token budget (
nemotron35_mtp6_t64), it answers… I know that the capital of China is Beijing.and the check isTrue— 13 decode graphs, 2 MTPlines. This is recorded rather than smoothed over because a token budget is not a serving result.
The one row that needed a base fix now has it. On the previous revision of this branch — whose
image was built from
5b8d95a4a9, i.e. before the base gained #260 — the MLA serving row failed:request 4, the first request large enough to build a chunked MLA context, killed the engine with
That is MUSA-100055 — an unfinished migration in the fork's own MLA reader, which upstream v0.28.0
re-laid-out from flat parallel lists to per-chunk objects. It was never attributable to these pins:
the reader is pure-Python fork code with no
torch/torchadasurface, the image's vLLM commit equalsthird_party/PINS, and the short-promptFLASHMLArow passed even then. It is fixed by #260, whichis merged into this base (
0fae3a1948;grep -c seq_toton the mergedcommon.pyreturns 0),and row 4 passes in the matrix above. So these pins are validated on the code that will actually
merge, not on a base that still carried the defect.
One log line to read carefully, because a reviewer grepping the r2 log will hit it: the passing row's
server log still contains one
EngineDeadError. It is at06:06:49, inside the shutdown window —[shutdown] Process manager: force killing remaining process EngineCore pid 7802→MPClient: engine manager stopped→ the API server'soutput_handlercoroutine, still awaiting engine output,raises — i.e. it is the runner's own SIGTERM tearing the engine down after the smoke, not a serving
failure. Contrast the failing r1 log, where the same line sits at
05:16:23mid-serving next toAttributeError: ... seq_totand an actualPOST /v1/chat/completions HTTP/1.1" 500. In r2 thereare 0 HTTP 500s, 5/5 requests return 200, and
seq_totoccurs 0 times in the server log (r1: 2).Not run, deliberately
deep_epis validated here by the image's build-time import gate and by theABI/library check above; no expert-parallel row was run, so the DeepEP all-to-all backend is
not exercised at runtime. Tracked as MUSA-100056.
Nemotron row is a single-request functional smoke on the compiled protocol, not a benchmark.
tests/beyond the two dependency-contract modules (not collectible in the review interpreter).Evidence
Build fingerprint, image push log, the full build log with the dependency gate, per-row server and
client logs, both probe payloads, the Nemotron re-run and the runner itself are attached to
MUSA-100057 in the workspace:
generated/MUSA-100057/build-fingerprint-r2.txt— branch, head, base, changed-file hashes, pinsgenerated/MUSA-100057/build-r2.log— rebuild including thePASS import …dependency gategenerated/MUSA-100057/matrix-run-r2/— r2 matrix:summary.txt,preflight.log, per-row logs(
qwen3_8b_bf16,qwen3_moe_fp8,dsv2lite_mla,dsv2lite_server+dsv2lite_client,nemotron35_mtp6),probe_like_device.log,probe_moe_config.log,runner-console.loggenerated/MUSA-100057/matrix-run-r2-nemotron-recheck/— the 64-token Nemotron re-rungenerated/MUSA-100057/matrix-run/— the previous revision's matrix, kept because it is the runthat exposed the MLA row that fix(mla): migrate MUSA MLA chunked-context prefill to the vLLM v0.28 chunk layout #260 then fixed