Repository navigation
[Perf] MiniMax-H3 VAE Decoder Ops - #6607
Conversation
| if HAS_TRITON: | ||
|
|
||
| @triton.jit | ||
| def _fp16_qk_rope_kernel( |
There was a problem hiding this comment.
model specific only?
There was a problem hiding this comment.
This PR is an experiment to explore the practical boundary of eager, model-specific optimization. I intentionally do not plan to move these kernels into a broad common operator layer.
In diffusion models, preserving model quality while continuously supporting new checkpoints, request policies, and execution features requires model-specific ownership. A “generic” operator still needs dtype, shape, device, stride, execution-mode, and platform checks. Even after those checks, it cannot automatically guarantee the model’s exact mathematical semantics, nor that every platform receives a positive performance gain. Those guarantees still have to be validated and maintained model by model.
Moving the implementation into a common layer would therefore not remove the model-specific work; it would mostly hide that work behind a wider API and create stronger coupling between models.
The approach in this PR keeps the boundary explicit: the model owns its operator contract, numerical policy, hardware dispatch, fallbacks, and validation. This makes the path easier to reason about, easier to dispatch safely, and easier for the community to validate or extend for another platform or quality level.
The trade-off is intentional duplication. Source-code volume alone is also becoming a progressively smaller maintenance constraint with AI-assisted tooling; But I’m open to other views.
There was a problem hiding this comment.
Fair position — with #6305 (op-boundary RFC) still settling where shared ops live, an explicitly model-owned kernel with its own contract, numerical policy, and validation is a defensible point on that spectrum. Could you capture this rationale in the PR description or a module docstring, so the placement reads as a decision rather than an oversight? If a future #6305 tier ever makes these shareable, migration can be a follow-up.
There was a problem hiding this comment.
Captured this in the model-local ops/README.md and the PR description in d34b75bd4. The README makes the numerical/execution/platform contract the reason for local ownership, documents the evidence expected for additional platforms, and leaves migration to a future #6305 shared tier as a follow-up. Thanks.
0d2054b to
8315dfa
Compare
Signed-off-by: mglyn <1203789601@qq.com>
Signed-off-by: mglyn <1203789601@qq.com>
Signed-off-by: mglyn <1203789601@qq.com>
Signed-off-by: mglyn <1203789601@qq.com>
|
Codex usage limits have been reached for code reviews. Please check with the admins of this repo to increase the limits by adding credits. |
|
This PR appears to be related to model: MinimaxH3. Model owners: @david6666666 @mglyn, 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. |
Signed-off-by: mglyn <1203789601@qq.com>
B300 validation (SM103)I tested base Important: the PR does not currently enable the optimization on B300
The B300 run of I also ran the complete official-shape VAE decode with seed 42, latent
So the current PR has no realized B300 performance gain, as expected from the fallback. SM103 feasibility check (forced dispatch, not current PR behavior)To evaluate the implementation itself, I forced the existing exact operator set to be selected on SM103 without changing any kernel code. This is a feasibility experiment only, not a claim that the PR currently supports B300. On the same loaded VAE and fixed latent, using two warmups followed by five CUDA-event measurements per side:
Correctness passed for the complete decoded tensor: Isolated production-shape operator measurements were also positive and bitwise exact:
I additionally ran 256 combinations of shape, seed, and value distribution for direct Q/K+RoPE and scaled-residual checks (512 bitwise comparisons total); all passed on SM103. Conclusion: the kernels appear both exact and beneficial on this B300, but the current hardware table prevents users from receiving the gain. B300 should get an explicitly registered and validated SM103 operator set (rather than being reported as supported by the current SM90-only dispatch). Once that is added, the five selected-operator tests should run instead of skip, and I recommend repeating the complete VAE A/B on the actual updated PR head. |
|
Ran a controlled local GPU A/B on the exact merge base
The one-time process-to-ready time increased from Environment: NVIDIA L20X (SM90), Python 3.12.13, PyTorch 2.13.0+cu129, CUDA 12.9, Triton 3.7.1, vLLM 0.28.0, MiniMax-H3 snapshot The tested path passed without a blocking correctness or integration issue. |
Signed-off-by: mglyn <1203789601@qq.com>
Signed-off-by: mglyn <1203789601@qq.com>
Signed-off-by: mglyn <1203789601@qq.com>
Signed-off-by: mglyn <1203789601@qq.com>
Signed-off-by: mglyn <1203789601@qq.com>
hsliuustc0106
left a comment
There was a problem hiding this comment.
Thanks for the rationale on the model-specific boundary — accepted, especially while #6305 is still settling where shared ops live. The one remaining ask from my side is to capture that intent in the PR description or a module docstring so the placement reads as a decision rather than an oversight.
|
This is already captured in the second paragraph of the PR description and in the model-local |
) vllm-project#6283's DiT SwiGLU fusion wiring stays reverted (transformer call sites removed); this generic layer file is still required by the H3 VAE decoder ops, which are shared by every bench group and are not part of the fusion A/B under test.
Signed-off-by: mglyn <1203789601@qq.com> Co-authored-by: Hongsheng Liu <liuhongsheng4@huawei.com> (cherry picked from commit 04b97a5)
Signed-off-by: mglyn <1203789601@qq.com> Co-authored-by: Hongsheng Liu <liuhongsheng4@huawei.com>
Signed-off-by: mglyn <1203789601@qq.com> Co-authored-by: Hongsheng Liu <liuhongsheng4@huawei.com>
This draft evaluates model-specific eager optimizations for the MiniMax-H3 video VAE decoder. The registered hardware entries target CUDA compute capabilities 9.0, 10.0, and 10.3, while hardware selection stays behind a flat operator-set dispatch so another backend can be added without changing the model integration. It is a follow-up and alternative optimization direction to #5979: instead of regionally compiling every VAE Transformer block, it accelerates the decoder's repeated eager operations.
The model-local placement is intentional: numerical correctness and performance depend on H3's exact math, dtype, layout, execution contract, and per-platform validation, so its dispatch and fallbacks remain beside the model integration. The adjacent
ops/README.mdrecords this boundary and the evidence expected when extending platform support. If #6305 establishes a suitable shared operator tier, migration can be a follow-up.Every optimized operator now preserves the reference output bit for bit. Q/K normalization and RoPE use one fused Triton kernel per Q/K tensor while retaining the reference RMS reduction order and explicit FP16 multiply/add rounding. Every eligible call uses the same exact path. Unsupported devices, tensor contracts, spatial-parallel attention, and compile execution retain the original module operations.
Kernel development followed MIT HAN Lab Kernel Design Agents, the evidence-first workflow from KDA-Pilot.
Operations
torch.equalfor Q and KSiluAndMulafter the FP16w1outputtorch.equalFP32 residual + FP16 branch * FP32 scale, with contraction disabledtorch.equalOnly decoder Transformer-block Linear weights are materialized in FP16. Keyframe encoding and non-block decoder parameters, including
proj_out, remain FP32.Complete VAE decode
The exact fused Q/K path was checked on one NVIDIA H200 with the official 1344x768 output shape, 209 output frames, latent shape
[1, 24, 62, 48, 84], and decoded tensor shape[1, 3, 209, 768, 1344]. The complete decoded tensor is bitwise equal to the reference Q/K implementation (torch.equal=True, max absolute error 0).The same operator set was independently exercised on B300/SM103 before dispatch registration: the complete decoded tensor remained bitwise equal, 512 direct exactness comparisons passed, and complete VAE decode P50 improved from 8.7199 s to 4.4145 s. The current head explicitly registers SM90, SM100, and SM103; other architectures retain the reference fallback.
Turbo end-to-end performance
The controlled A/B used the MiniMax-H3 Turbo v1.0 FL2VA recipe on NVIDIA H200 GPUs: 768x1344 (
width x height), 107 frames at 24 FPS, 4 denoiser evaluations, regional DiT compile, and the same prompt, image, and seed. The single-GPU case used USP1 and VAE tile PP1; the four-GPU case used USP4, ring 1, and VAE tile PP4. Both commits received two complete warmup requests followed by three measured requests. Stage 0 comes from the online benchmark response. VAE decode comes from the pipeline profiler; the four-GPU value is the per-request critical path, or the maximum across ranks.Diffuse P50 remained effectively unchanged: 21.578 s versus 21.593 s on one GPU and 5.960 s versus 5.964 s on four GPUs. This isolates the end-to-end improvement to VAE decode. The baseline is main commit
c6df39e9f; the single-GPU candidate isf67759666, and the four-GPU candidate is6030e5976.Compile execution bypasses these operators and retains the original expressions for Inductor.
Integration and fallback
Validation
The final validation was repeated after rebasing onto main commit
c6df39e9f.git diff --checktests/diffusion/models/minimax_h3, excluding the external full-model caseThe operator tests cover production and regression token lengths, exact outputs, unsupported-input fallback, hardware dispatch, idempotent installation, official-semantics guards, incompatible remote-code structure, selective Linear conversion, and the post-Linear FP16 SwiGLU path.