Repository navigation
feat(magi2): add fused SwiGLU7 activation - #7252
Conversation
9f114b0 to
61c943d
Compare
01c53c5 to
028f8ec
Compare
|
This PR appears to belong to: docs/design/module/benchmarking.md, docs/design/module/diffusion/diffusion_model_integration.md, docs/design/module/diffusion/index.md. Module owners: @wtomin @RuixiangMa @david6666666 Routing: @wtomin via module of the changed files, module named in the PR description, semantic router, CODEOWNERS; @RuixiangMa via module of the changed files, module named in the PR description, semantic router; @david6666666 via module of the changed files, module named in the PR description @yeahdongcn, please review your own changes and leave a short self-review comment describing what you checked. PRs without author self-review may not be assigned a reviewer. Please take a look when you have a chance. If you would like an automated review, mention @vllm-omni-review-bot in a comment. |
hsliuustc0106
left a comment
There was a problem hiding this comment.
The kernel, fallback matrix, and tests all verify clean on my side (19 CPU + 14 CUDA tests pass; parity exact on CPU and within ULP-scale tolerances on GPU). The only blocker is the unregistered MAGI2_USE_FUSED_SWIGLU7 env var — the environment-inventory drift test fails at this head, so the cpu lane would go red on merge.
Independently reproduced the perf claim on L20X (triton 3.7.1): I=1280 gives 1.94× (256 tok) to 3.31× (4096 tok), and the dense-MLP width I=12288 reaches 8.08× at 4096 tokens with parity passing — the reported 1.7× understates the gain at production widths. Approve-able once the registry entry lands.
477b240 to
2a2eceb
Compare
Omni ReviewBot: no human activity for 7 days@yeahdongcn this pull request has had no human commit, comment or review since 2026-09-16. Please confirm the current plan and next step. The author or a maintainer decides whether to change the PR state. To keep it moving, any one of these is enough: push an update, reply to the open blocker, or post the current plan and timeline. |
| linear = tl.load(x_ptr + 2 * offsets + 1, mask=mask, other=0).to(tl.float32) | ||
| # Comparisons preserve NaNs; gate is only upper-clamped, unlike linear. | ||
| gate = tl.where(gate > 7.0, 7.0, gate) | ||
| linear = tl.where(linear < -7.0, -7.0, linear) |
There was a problem hiding this comment.
why not use a general swiglu with 7 as an input
hsliuustc0106
left a comment
There was a problem hiding this comment.
Env-var blocker resolved by removing the flag entirely, and the tests now run in the core_model tier. Approving; the general-SwiGLU-shape question above stands as a non-blocking follow-up.
Omni ReviewBot: supersededThe CI failure noted on |
Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>
2a2eceb to
3aadf82
Compare
Summary
Add MAGI-2's packed SwiGLU7 activation as a diffusion
CustomOpand use thefused Triton implementation by default for eager GPU execution. CPU execution
and
torch.compilekeep the native PyTorch expression.The implementation preserves the released formula: FP32 arithmetic, upper-only
gate clamp, symmetric linear clamp,
linear + 1, and one final output cast.It uses no MUSA-only API and does not require an environment variable.
Validation
Results on commit
2a2eceb93:10.20.36.18, GPU 7 (48 SM) / MUSA / Triton 3.2.0: 14 passedMicrobenchmark:
10.20.36.18GPU 7 (48 SM), Triton 3.2.0Both runs passed numerical parity. The benchmark uses alternating native/fused
calls with a 512 MiB cache flush outside timed regions; it is an eager operator
diagnostic rather than a full-video throughput claim.
Scope
No MUSA-specific namespace patch, MATE dependency, mHC, MoE, EP, sampler, or
platform-wide change is included. Independent of #7206, #7249, and #7261;
#7156 remains the integration reference.