[Performance] Performance improvements in non-blockwise fp8 CUTLASS MoE #20762
Conversation
Signed-off-by: ElizaWszola <ewszola@redhat.com>
|
👋 Hi! Thank you for contributing to the vLLM project. 💬 Join our developer Slack at https://slack.vllm.ai to discuss your PR in #pr-reviews, coordinate on features in #feat- channels, or join special interest groups in #sig- channels. Just a reminder: PRs would not trigger full CI run by default. Instead, it would only run Once the PR is approved and ready to go, your PR reviewer(s) can run CI to test the changes comprehensively before merging. To run CI, PR reviewers can either: Add 🚀 |
|
This pull request has merge conflicts that must be resolved before it can be |
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
There was a problem hiding this comment.
Summary of Changes
Hello @ElizaWszola, 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 significantly enhances the performance of non-blockwise FP8 CUTLASS Mixture-of-Experts (MoE) operations. It achieves this by re-integrating and optimizing the pre-calculation of GEMM strides and by deploying more efficient CUDA kernels for data shuffling within the MoE layer, leading to overall speed improvements.
Highlights
- Re-enabled Stride Pre-calculation: The
ab_stridesandc_stridesparameters, crucial for efficient GEMM operations in CUTLASS MoE, are now pre-calculated and passed explicitly to the kernels. This re-enables a performance optimization that was previously disabled due to PPLX compatibility issues, which are now resolved. - Optimized Data Shuffling Kernels: Introduced and integrated faster CUDA kernels for shuffling input hidden states, their scales, and the final outputs of the MoE layer. A fallback 'slow' kernel is also provided for cases where column alignment prevents the use of the fastest kernel, ensuring robustness.
- API and Architecture Updates: The
run_cutlass_moe_fp8function and theCutlassExpertsFp8class have been updated to acceptab_stridesandc_stridesas explicit arguments. The responsibility for calculating these strides has been centralized withinCompressedTensorsW8A8FP8MoEMethod.process_weights_after_loading, streamlining their management.
Using Gemini Code Assist
The 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 in your pull request via creating an issue comment (i.e. comment on the pull request page) using either /gemini <command> or @gemini-code-assist <command>. Below is a summary of the supported commands.
| Feature | Command | Description |
|---|---|---|
| Code Review | /gemini review |
Performs a code review for the current pull request in its current state. |
| Pull Request Summary | /gemini summary |
Provides a summary of the current pull request in its current state. |
| Comment | @gemini-code-assist | Responds in comments when explicitly tagged, both in issue comments and review comments. |
| Help | /gemini help |
Displays a list of available commands. |
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 .gemini/ folder in the base of the repository. Detailed instructions can be found here.
Limitations & Feedback
Gemini Code Assist is currently in preview and 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 to provide feedback.
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
-
Review the Privacy Notices, Generative AI Prohibited Use Policy, Terms of Service, and learn how to configure Gemini Code Assist in GitHub here. Gemini can make mistakes, so double check it and use code with caution. ↩
There was a problem hiding this comment.
Code Review
This pull request introduces performance improvements for non-blockwise fp8 CUTLASS MoE. The main changes include pre-calculating strides and using faster shuffle kernels. The changes are well-contained and affect benchmark, test, and kernel implementation files.
My review identified a potential high-severity bug in vllm/model_executor/layers/quantization/compressed_tensors/compressed_tensors_moe.py where a function call is missing required arguments, which could lead to a runtime error. I also found a medium-severity maintainability issue in csrc/moe/moe_permute_unpermute_op.cu related to code duplication. I've provided suggestions to fix both issues. Overall, the changes align with the stated performance goals.
| if self.fused_experts is None: | ||
| # If no modular kernel is provided, use cutlass_moe_fp8 | ||
| from vllm.model_executor.layers.fused_moe.cutlass_moe import ( | ||
| cutlass_moe_fp8) | ||
| return cutlass_moe_fp8( | ||
| x, | ||
| layer.w13_weight, | ||
| layer.w2_weight, | ||
| topk_weights, | ||
| topk_ids, | ||
| activation=activation, | ||
| global_num_experts=global_num_experts, | ||
| expert_map=None if self.disable_expert_map else expert_map, | ||
| w1_scale=layer.w13_weight_scale, | ||
| w2_scale=layer.w2_weight_scale, | ||
| ab_strides1=self.ab_strides1, | ||
| ab_strides2=self.ab_strides2, | ||
| c_strides1=self.c_strides1, | ||
| c_strides2=self.c_strides2, | ||
| a1_scale=layer.w13_input_scale, | ||
| a2_scale=layer.w2_input_scale, | ||
| ) |
There was a problem hiding this comment.
The call to cutlass_moe_fp8 in the if self.fused_experts is None: block is missing the new stride arguments (ab_strides1, ab_strides2, c_strides1, c_strides2). Since the function signatures for MoE kernels are being updated in this PR to require these strides, this will likely cause a TypeError if this code path is executed. Please pass the stride tensors to the cutlass_moe_fp8 call.
return cutlass_moe_fp8(
x,
layer.w13_weight,
layer.w2_weight,
topk_weights,
topk_ids,
activation=activation,
global_num_experts=global_num_experts,
expert_map=None if self.disable_expert_map else expert_map,
w1_scale=layer.w13_weight_scale,
w2_scale=layer.w2_weight_scale,
ab_strides1=self.ab_strides1,
ab_strides2=self.ab_strides2,
c_strides1=self.c_strides1,
c_strides2=self.c_strides2,
a1_scale=layer.w13_input_scale,
a2_scale=layer.w2_input_scale,
)| if (num_cols % (128 / sizeof(input_tensor.scalar_type()) / 8)) { | ||
| // use slow kernel if num_cols can't be aligned to 128 bits | ||
| MOE_DISPATCH(input_tensor.scalar_type(), [&] { | ||
| shuffleInputRowsKernelSlow<scalar_t><<<blocks, threads, 0, stream>>>( | ||
| reinterpret_cast<scalar_t*>(input_tensor.data_ptr()), | ||
| dst2src_map.data_ptr<int32_t>(), | ||
| reinterpret_cast<scalar_t*>(output_tensor.data_ptr()), num_src_rows, | ||
| num_dest_rows, num_cols); | ||
| }); | ||
| } else { | ||
| MOE_DISPATCH(input_tensor.scalar_type(), [&] { | ||
| shuffleInputRowsKernel<scalar_t><<<blocks, threads, 0, stream>>>( | ||
| reinterpret_cast<scalar_t*>(input_tensor.data_ptr()), | ||
| dst2src_map.data_ptr<int32_t>(), | ||
| reinterpret_cast<scalar_t*>(output_tensor.data_ptr()), num_src_rows, | ||
| num_dest_rows, num_cols); | ||
| }); | ||
| } |
There was a problem hiding this comment.
The MOE_DISPATCH macro is duplicated in both the if and else branches. This can be refactored to have a single MOE_DISPATCH call with the conditional logic inside the lambda to improve code readability and maintainability by reducing duplication.
MOE_DISPATCH(input_tensor.scalar_type(), [&] {
if (num_cols % (128 / sizeof(input_tensor.scalar_type()) / 8)) {
// use slow kernel if num_cols can't be aligned to 128 bits
shuffleInputRowsKernelSlow<scalar_t><<<blocks, threads, 0, stream>>>(
reinterpret_cast<scalar_t*>(input_tensor.data_ptr()),
dst2src_map.data_ptr<int32_t>(),
reinterpret_cast<scalar_t*>(output_tensor.data_ptr()), num_src_rows,
num_dest_rows, num_cols);
} else {
shuffleInputRowsKernel<scalar_t><<<blocks, threads, 0, stream>>>(
reinterpret_cast<scalar_t*>(input_tensor.data_ptr()),
dst2src_map.data_ptr<int32_t>(),
reinterpret_cast<scalar_t*>(output_tensor.data_ptr()), num_src_rows,
num_dest_rows, num_cols);
}
});
Signed-off-by: ElizaWszola <ewszola@redhat.com>
|
This pull request has merge conflicts that must be resolved before it can be |
Signed-off-by: ElizaWszola <ewszola@redhat.com>
Signed-off-by: ElizaWszola <ewszola@redhat.com>
ElizaWszola
requested review from
tlrmchlsmth,
WoosukKwon,
mgoin and
robertgshaw2-redhat
as code owners
Signed-off-by: ElizaWszola <ewszola@redhat.com>
|
This pull request has merge conflicts that must be resolved before it can be |
This PR introduces a couple performance improvements to non-blockwise fp8 CUTLASS MoE. The improvements are:
ab_stridesandc_strides. It had been disabled due to PPLX codepath compatibility issues which are now resolved.hidden_states, input scales and outputs of the function that executes CUTLASS MoE layer.Benchmarks
Execution times of function
cutlass_moe_fp8times in microseconds (μs):