Conversation
Signed-off-by: Zhen Xie <zxie3@binghamton.edu>
Signed-off-by: Zhen Xie <zxie3@binghamton.edu>
Signed-off-by: Zhen Xie <zxie3@binghamton.edu>
Signed-off-by: Zhen Xie <zxie3@binghamton.edu>
zhenhuaw-me
left a comment
There was a problem hiding this comment.
Thank you Zhen Xie for the PR! I am very happy to see community contribution. Could you please follow https://github.com/NVIDIA/TensorRT-LLM/blob/main/CONTRIBUTING.md and bring this PR to product support shape?
| int const seqlen_per_bs = static_cast<int>(query.size(1)); | ||
| // Scheduling-only tuning knob. All supported values execute one warp per | ||
| // head and preserve identical arithmetic/order within that warp. | ||
| static int const blockSize = [] { |
There was a problem hiding this comment.
Since the kernel is implemented for Wan, does it make sense to drop the env?
|
Hi Zhen, seems that you've contributed to the wrong branch |
luyiyun1021
left a comment
There was a problem hiding this comment.
Typically, we would both need kernel unittests and e2e lpips comparison to ensure there is no accuracy regression. check scripts/visualgen_eval/visual_gen_lpips_score_eval.py
| sumOfSquares += vals.x * vals.x + vals.y * vals.y; | ||
| } | ||
| sumOfSquares = llm::common::warpReduceSum(sumOfSquares); | ||
| float const rms_rcp = rsqrtf(sumOfSquares / static_cast<float>(head_dim) + eps); |
There was a problem hiding this comment.
Wan's norm_q/norm_k normalize over the full hidden dimension (num_heads * head_dim, i.e. 5120 for A14B), before splitting into heads. Here each warp only accumulates one head's 128 elements, which changes the operation to per-head RMSNorm. Could we reduce across all heads of each token and add a parity test with different input magnitudes across heads?
| // Match the unfused path's BF16 RMSNorm output before the FP32 RoPE | ||
| // multiply instead of carrying extra precision across the fusion. | ||
| elements[i] = __bfloat162float(__float2bfloat16_rn( | ||
| elements[i] * rms_rcp * __bfloat162float(weight[dim]))); |
There was a problem hiding this comment.
dim is local to the current head, but Wan's RMSNorm weights span all num_heads * head_dim dimensions. This currently reuses the first head's weights for every head. Should this index be weight[headIdx * head_dim + dim]? A parity test with nonuniform norm weights across heads would catch this independently of the reduction-domain issue.
| offsets = tl.program_id(0) * block + tl.arange(0, block) | ||
| mask = offsets < count | ||
| values = tl.load(x + offsets, mask=mask, other=0.0).to(tl.float32) | ||
| output = (values * tl.sigmoid(1.702 * values)).to(tl.bfloat16) |
There was a problem hiding this comment.
Wan's FFN uses activation_fn="gelu-approximate", which resolves to GELU(approximate="tanh") in the pinned diffusers 0.36.0. This helper computes x * sigmoid(1.702 * x), a different activation. The .proj guard also accepts the original GELU module, so the optimized path replaces Wan's activation. Could we fuse the tanh GELU formula and validate parity against the actual Wan FFN? The dispatch should also check the activation type and approximation setting.
o-stoner
left a comment
There was a problem hiding this comment.
Thanks for the writeup! For change 1, this looks like it may already be implemented on main; see Attention.forward() / apply_packed_qk_norm_rope() in tensorrt_llm/_torch/visual_gen/modules/attention.py (wired into Wan via fuse_qk_norm_rope in transformer_wan.py). Could you check whether your changes should be perhaps used to extend the existing kernel instead?
[visual_gen] Optimize Wan2.2 Dataflow (Zhen Xie from VibeHPC)
Description
This PR adds opt-in dataflow optimizations for Wan 2.2 inference in
visual_gen, targeting single-GPU execution on NVIDIA B300.The changes reduce QKV packing overhead, eliminate redundant FP8 activation quantization around the FFN, cache immutable FP8 weight metadata, and tune the fused Q/K normalization and RoPE kernel launch configuration.
All optimizations are disabled by default. This PR does not modify the model architecture, sampling procedure, MLPerf harness, dataset, request scheduler, or output format.
Changes
1. Preserve Packed QKV Storage Through Q/K Normalization and RoPE
The fused QKV projection already produces Q, K, and V in one shared allocation.
The original Q/K normalization and RoPE path creates separate output tensors. As a result, Transformer Engine repacks Q, K, and V before FP8 attention by launching a
CatArrayBatchedCopykernel.This PR adds a Wan-specific CUDA kernel that:
[Q | K | V]storage layout.Transformer Engine can then recognize the inputs as packed
bsh3dstorage and skip its per-attentionCatArrayBatchedCopy.The fused kernel preserves the BF16 rounding point between RMSNorm and RoPE to match the original Wan computation order.
Relevant files:
visual_gen/csrc/DiTRMSNormRope/fused_qk_norm_rope_kernel.cuvisual_gen/models/transformers/wan_transformer.py2. Fuse Wan ApproximateGELU With FP8 Quantization
Wan uses the following FFN activation:
The original execution path is:
This PR adds a Triton implementation that:
Float8Tensor.The optimized execution path is:
The FFN-down projection consumes the generated
Float8Tensordirectly, avoiding a second activation amax and FP8 cast.Unsupported tensor layouts and non-inference execution continue to use the original FeedForward path.
Relevant files:
visual_gen/ops/wan_fused.pyvisual_gen/models/transformers/wan_transformer.pyvisual_gen/ops/linear.py3. Cache Immutable FP8 Weight Wrappers
The FP8 linear path previously reconstructed the following objects during repeated inference calls:
Float8Tensorwrapper;Model weights remain unchanged during inference, so recreating these objects introduces unnecessary overhead.
This PR caches the prepared FP8 weight representation using:
The cache automatically refreshes if either the weight or its scale changes.
Pointer and version checks execute outside Torch Dynamo tracing to avoid graph breaks and
DataPtrVariableerrors.Relevant file:
visual_gen/ops/linear.py4. Tune the Wan Q/K Norm and RoPE Launch Configuration
The fused Wan Q/K Norm and RoPE kernel supports the following thread-block sizes:
This PR changes the default to 512 threads for the validated Wan 2.2 configuration:
Each warp continues to process one Q or K head row. The optimization only changes block scheduling and does not change the arithmetic performed within a warp.
The launch configuration can be overridden with:
Relevant file:
visual_gen/csrc/DiTRMSNormRope/fused_qk_norm_rope_kernel.cuPerformance
Measurements were collected on a single NVIDIA B300 GPU using the Wan 2.2 A14B MLPerf Offline workload.
Test Configuration
Isolated Operator Results
Packed QKV Profiling Details
The real attention shape is:
The original path contains one QKV packing kernel per FP8 attention call:
Measured results:
Multiple Nsight Systems measurements indicated a local saving of 1.4 ms per attention invocation, depending on the run.
ApproximateGELU and FP8 Quantization Profiling Details
The real FFN activation shape is:
Three independent long-run measurements produced:
The selected Triton quantization block size is:
End-to-End Performance Comparison
The weight-cache result is close to normal run-to-run variation. The optimization remains included because it removes redundant preparation work and did not introduce an observed memory or correctness issue.
Q/K Norm and RoPE Thread Sweep
The 512-thread configuration measured 0.21% higher QPS than the 256-thread configuration.
Performance Result Scope
These measurements are development A/B results.
They are not MLPerf submission-valid performance results because the short exploratory runs do not satisfy the full MLPerf duration and query-count requirements.
Numerical Behavior
Packed QKV
Packed QKV and separate QKV produced identical outputs for the isolated attention comparison:
The optimized path retains the same cuDNN FP8 SDPA kernel:
The optimization changes the QKV storage layout presented to Transformer Engine but does not replace the attention computation.
Fused ApproximateGELU and FP8 Quantization
The fused kernel preserves Wan's activation formula:
Isolated numerical checks produced:
Fused Q/K Norm and RoPE
The original path rounds the RMSNorm output to BF16 before applying the FP32 RoPE multiplication.
The fused CUDA kernel explicitly preserves this rounding boundary:
This prevents the fusion from unintentionally carrying additional FP32 precision across the original BF16 boundary.
Accuracy Validation
The isolated numerical tests passed for the optimized operator paths.
Full MLPerf
AccuracyOnlyand VBench validation should still be completed before enabling these optimizations in a submission configuration.Enabling the Optimizations
All behavior-changing paths are disabled by default.
Baseline
unset VISUAL_GEN_WAN_OPTIMIZATIONSEnable the Optimized Wan Path
When the master switch is enabled, the following optimizations are enabled by default:
Disable Individual Components
Disable Packed QKV:
export VISUAL_GEN_WAN_PACKED_QKV=0Disable fused ApproximateGELU and FP8 quantization:
export VISUAL_GEN_WAN_FUSED_GELU_QUANT=0Disable the FP8 weight cache:
export VISUAL_GEN_WAN_FP8_WEIGHT_CACHE=0Use the original 256-thread Q/K Norm and RoPE configuration:
export VISUAL_GEN_WAN_QK_ROPE_THREADS=256The worker process must be restarted after changing these variables because the configuration is read when the Python modules are imported.
Fallback Behavior
The optimized FFN path checks the following conditions before using the fused ApproximateGELU and FP8 quantization implementation:
If any condition is not satisfied, execution falls back to the original FeedForward implementation.
The packed QKV path checks:
Unsupported cases use the original Q/K normalization and RoPE path.
Scope
This PR only changes the Wan 2.2 inference implementation in
visual_gen.It does not change:
Files Changed
visual_gen/csrc/DiTRMSNormRope/fused_qk_norm_rope_kernel.cuAdds:
visual_gen/models/transformers/wan_transformer.pyAdds:
visual_gen/ops/wan_fused.pyAdds:
Float8Tensorconstruction.visual_gen/ops/linear.pyAdds:
Float8Tensor;visual_gen/README-for-dataflow-optimizations.mdDocuments:
Summary
This PR improves Wan 2.2 inference dataflow without changing the model architecture or benchmark framework.
The main improvements are: