[diffusion]: optimize zimage with rmsnorm + add + (rmsnorm)#16404
[diffusion]: optimize zimage with rmsnorm + add + (rmsnorm)#16404attack204 wants to merge 1 commit intosgl-project:mainfrom
Conversation
attack204
requested review from
BBuf,
DarkSharpness,
mickqian and
yhyang201
as code owners
Summary of ChangesHello @attack204, 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 pull request focuses on optimizing the Highlights
🧠 New Feature in Public Preview: You can now enable Memory 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 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 counter productive. 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. You can also get AI-powered code generation, chat, as well as code reviews directly in the IDE at no cost with the Gemini Code Assist IDE Extension. Footnotes
|
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
There was a problem hiding this comment.
Code Review
This pull request introduces optimized fused kernels for RMSNorm operations within the ZImage model, which yields a significant performance improvement. The changes are well-structured, including new CUDA kernels, Python wrappers for JIT compilation, and seamless integration into the model with a fallback mechanism for compatibility. My review has identified one critical issue that would cause a runtime error, along with a few medium-severity suggestions to enhance code maintainability, robustness, and performance. Addressing these points will further solidify the quality of this optimization.
| output1: intermediate result (x + RMSNorm(residual)) [batch_size, seq_len, hidden_size] | ||
| output2: final result (RMSNorm(output1)) [batch_size, seq_len, hidden_size] | ||
| """ | ||
| can_use_cuda = _is_cuda and (x.shape[-1] % 4 == 0) |
There was a problem hiding this comment.
The variable _is_cuda is not defined, which will cause a NameError at runtime. It seems you intended to check if the tensor is on a CUDA device. You should use x.is_cuda instead. This issue is also present on line 583 for the FusedRMSNormAdd class.
| can_use_cuda = _is_cuda and (x.shape[-1] % 4 == 0) | |
| can_use_cuda = x.is_cuda and (x.shape[-1] % 4 == 0) |
| # TODO: workaround, do not import cutlass from flashinfer | ||
| cutlass_include = os.path.join( | ||
| os.path.dirname(flashinfer.__file__), "data", "cutlass", "include" | ||
| ) |
There was a problem hiding this comment.
The cutlass_include path is constructed based on the internal directory structure of the flashinfer package. This creates a brittle dependency that might break if flashinfer changes its packaging in a future version. As noted by the TODO, this is a workaround. For long-term stability, it would be more robust to find a way to get this path through flashinfer's public API if available, or consider vendoring the required headers. This comment also applies to python/sglang/jit_kernel/diffusion/fused_rmsnorm_add_rmsnorm.py.
| def _rmsnorm(self, x: torch.Tensor, weight: Optional[torch.Tensor]) -> torch.Tensor: | ||
| """Fallback RMSNorm implementation""" | ||
| variance = x.float().pow(2).mean(dim=-1, keepdim=True) | ||
| x_normalized = x * torch.rsqrt(variance + self.eps) | ||
| if weight is not None: | ||
| x_normalized = x_normalized * weight | ||
| return x_normalized.to(x.dtype) |
There was a problem hiding this comment.
The _rmsnorm method is duplicated in both FusedRMSNormAddRMSNorm (lines 541-547) and FusedRMSNormAdd (lines 602-608). To improve maintainability and reduce code duplication, consider refactoring this into a static method or a standalone helper function.
|
Do we have any test/benchmark on this kernel in |
| #include <tvm/ffi/container/tensor.h> | ||
| #include <tvm/ffi/optional.h> | ||
|
|
||
| #include "cutlass/numeric_types.h" |
There was a problem hiding this comment.
Can we try to eliminate this dependency? It seems only numeric types are used.
|
After this MR if (dtype.code == kDLFloat && dtype.bits == 32) {
dispatch_ipt(DTypeTag<float4, float>{});
} else if (dtype.code == kDLFloat && dtype.bits == 16) {
dispatch_ipt(DTypeTag<half4, half>{});
} else if (dtype.code == kDLBfloat && dtype.bits == 16) {
dispatch_ipt(DTypeTag<bf16_4, cutlass::bfloat16_t>{});
} |
|
What's the relationship of this PR and #14717? It seems there's many redundant code between these 2 PRs. We may need to maximize code reuse and avoid too many similar code. |
We have a kernel fusion plan for diffusion kernel. Many of which involve I summarized how we currently end up generating a large number of kernel source variants:
|
This PR is based on 14717 and uses some of its utilities functions, so after 14717 is merged, this PR will be rebased. |
4 participants
E2E IMPROMENT
10s/13% faster
python python/sglang/multimodal_gen/benchmarks/compare_perf.py baseline.json new.json ### Performance Comparison Report1. High-level Summary
2. Stage Breakdown
DETAIL PERF
TODO