-
-
Notifications
You must be signed in to change notification settings - Fork 7.2k
[Perf]Optimize rotary_emb implementation to use Triton operator for improved inference performance #16457
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Conversation
👋 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 🚀 |
a115f1f
to
e630289
Compare
Perhaps using kernel from https://github.com/vllm-project/flash-attention would be more reasonable |
39b1672
to
b2379fd
Compare
c435afc
to
8d6303f
Compare
Shouldn't this triton kernel have been used for Qwen2-VL? vllm/vllm/model_executor/models/qwen2_vl.py Lines 233 to 243 in 87e067d
vllm/vllm/model_executor/models/qwen2_5_vl.py Lines 276 to 282 in 87e067d
|
qwen2_vl.py
qwen2.py
@Isotr0py In qwen2-vl, the Qwen2Model of qwen2 will be called, and the rotary-embedding in Qwen2Model does not use the Triton kernel. During the token ID generation stage, there will be a performance bottleneck, which accounts for 40-60% of the overall inference time |
if use_flash_attn: | ||
return apply_rotary_emb(x.unsqueeze(0), cos, sin, | ||
not is_neox_style).squeeze(0) | ||
else: | ||
return _apply_rotary_emb_torch(x, cos, sin, is_neox_style) |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
if use_flash_attn: | |
return apply_rotary_emb(x.unsqueeze(0), cos, sin, | |
not is_neox_style).squeeze(0) | |
else: | |
return _apply_rotary_emb_torch(x, cos, sin, is_neox_style) | |
if current_platform.is_cuda_alike(): | |
from vllm_flash_attn.layers.rotary import apply_rotary_emb | |
return apply_rotary_emb(x.unsqueeze(0), cos, sin, | |
not is_neox_style).squeeze(0) | |
else: | |
return _apply_rotary_emb_torch(x, cos, sin, is_neox_style) |
Since triton is only available for Nvidia and ROCm GPUs, we can simplify the implementation here.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
hi,isotr0py,if I replace
from flash_attn.layers.rotary import apply_rotary_emb
return apply_rotary_emb(x.unsqueeze(0), cos, sin,
not is_neox_style).squeeze(0)
to
from vllm_flash_attn.layers.rotary import apply_rotary_emb
return apply_rotary_emb(x.unsqueeze(0), cos, sin,
not is_neox_style).squeeze(0)
Running the CI test will result in an error:
https://buildkite.com/vllm/ci/builds/18001#01964d1c-6acb-48b6-8b8e-6b296cdefd17
[2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] File "/usr/local/lib/python3.12/dist-packages/vllm/model_executor/models/qwen2.py", line 243, in forward
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] hidden_states = self.self_attn(
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] ^^^^^^^^^^^^^^^
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] File "/usr/local/lib/python3.12/dist-packages/torch/nn/modules/module.py", line 1739, in _wrapped_call_impl
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] return self._call_impl(*args, **kwargs)
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] File "/usr/local/lib/python3.12/dist-packages/torch/nn/modules/module.py", line 1750, in _call_impl
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] return forward_call(*args, **kwargs)
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] File "/usr/local/lib/python3.12/dist-packages/vllm/model_executor/models/qwen2.py", line 176, in forward
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] q, k = self.rotary_emb(positions, q, k)
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] File "/usr/local/lib/python3.12/dist-packages/torch/nn/modules/module.py", line 1739, in _wrapped_call_impl
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] return self._call_impl(*args, **kwargs)
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] File "/usr/local/lib/python3.12/dist-packages/torch/nn/modules/module.py", line 1750, in _call_impl
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] return forward_call(*args, **kwargs)
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] File "/usr/local/lib/python3.12/dist-packages/vllm/model_executor/layers/rotary_embedding.py", line 992, in forward
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] query_rot = _apply_rotary_emb(query_rot, cos, sin, self.is_neox_style)
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] File "/usr/local/lib/python3.12/dist-packages/vllm/model_executor/layers/rotary_embedding.py", line 81, in _apply_rotary_emb
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] from vllm_flash_attn.layers.rotary import apply_rotary_emb
| [2025-04-19T08:53:45Z] ERROR 04-19 01:53:45 [core.py:386] ModuleNotFoundError: No module named 'vllm_flash_attn'
How should I solve it
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I use flash.attn.layers.gross import
Mainly because qwen2-vl.py is used in this way, as follows
def apply_rotary_pos_emb_vision(t: torch.Tensor,
freqs: torch.Tensor,
use_flash_attn=False) -> torch.Tensor:
t_ = t.float()
cos = freqs.cos()
sin = freqs.sin()
apply_rotary_emb = apply_rotary_emb_torch
if use_flash_attn:
from flash_attn.layers.rotary import apply_rotary_emb
output = apply_rotary_emb(t_, cos, sin).type_as(t)
return output
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Oh, the import code should be vllm.vllm_flash_attn.layers.rotary import apply_rotary_emb
, otherwise it won't work outside the repo folder.
Qwen2-vl use original FA due to vllm_flash_attn
compatibility issue for ViT, however, we shouldn't use it here since original FA is not a requirement.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Using vllm-vllm_flash_mattn.layers.total still encounters errors:
https://buildkite.com/vllm/ci/builds/18013#01964e8f-a7a0-4638-baa0-a57dfa5ac8b4
| [2025-04-19T15:42:15Z] ERROR 04-19 08:42:15 [core.py:386] return forward_call(*args, **kwargs)
| [2025-04-19T15:42:15Z] ERROR 04-19 08:42:15 [core.py:386] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
| [2025-04-19T15:42:15Z] ERROR 04-19 08:42:15 [core.py:386] File "/usr/local/lib/python3.12/dist-packages/vllm/model_executor/layers/rotary_embedding.py", line 992, in forward
| [2025-04-19T15:42:15Z] ERROR 04-19 08:42:15 [core.py:386] query_rot = _apply_rotary_emb(query_rot, cos, sin, self.is_neox_style)
| [2025-04-19T15:42:15Z] ERROR 04-19 08:42:15 [core.py:386] ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
| [2025-04-19T15:42:15Z] ERROR 04-19 08:42:15 [core.py:386] File "/usr/local/lib/python3.12/dist-packages/vllm/model_executor/layers/rotary_embedding.py", line 81, in _apply_rotary_emb
| [2025-04-19T15:42:15Z] ERROR 04-19 08:42:15 [core.py:386] from vllm.vllm_flash_attn.layers.rotary import apply_rotary_emb
| [2025-04-19T15:42:15Z] ERROR 04-19 08:42:15 [core.py:386] ModuleNotFoundError: No module named 'vllm.vllm_flash_attn.layers'
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Are you saying that I need to make modifications https://github.com/vllm-project/flash-attention project ?
Can I modify it like this:
- Create a 'layers' sub directory in the' vllm_flash-attn 'directory
- Then copy 'flash_attn/layers/total. py' to the 'vllm_flash-attn/layers/' directory,
- Import or export 'apply_rotary-emb' in 'vllm_flash-attn/layers/init. py'.
- Ensure that the top-level 'init. py' of 'vllm_flash-attn' correctly exports the 'layers' module.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Are you saying that I need to make modifications https://github.com/vllm-project/flash-attention project ?
Yes. We should not directly copy the source code from FA to vllm_flash_attn
, because the copying should happen at the compilation during installation. So I recommend to make modifications in https://github.com/vllm-project/flash-attention to make rotary code copied during compilation.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
The file I want to copy/flash attention/flash attn/layers/rotary.py
contains the following code
from flash_attn.ops.triton.rotary import apply_rotary
If I copy /flash-attention/flash-attn/layers/rotary.py
to the /flash-attention/vllm_flash-attn/
directory during compilation, the error will still occur during execution,
I need to modify the code to become
from vllm.vllm_flash_attn.ops.triton.rotary import apply_rotary
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Are you saying that I need to make modifications https://github.com/vllm-project/flash-attention project ?
Yes. We should not directly copy the source code from FA to
vllm_flash_attn
, because the copying should happen at the compilation during installation. So I recommend to make modifications in https://github.com/vllm-project/flash-attention to make rotary code copied during compilation.
Can you help me merge this flash attn PR?
vllm-project/flash-attention#64
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
@Isotr0py hi,I have merged the changes to the flash attn repository. Can you help me review and merge this PR again
4dca833
to
679db29
Compare
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Overall LGTM now, just leave some nits. PTAL!
73354b5
to
bf7c6fc
Compare
Let's hold this off before vllm-project/flash-attention#64 merge.
This pull request has merge conflicts that must be resolved before it can be |
0069947
to
c04e840
Compare
…mproved inference performance Signed-off-by: cynthieye <[email protected]> Co-authored-by: MagnetoWang <[email protected]>
This pull request has merge conflicts that must be resolved before it can be |
…mproved inference performance Signed-off-by: cynthieye <[email protected]> Co-authored-by: MagnetoWang <[email protected]>
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM now, thanks for your patience!
* [BugFix] Remove default multiproc executor `collective_rpc` timeout (vllm-project#17000) Signed-off-by: Nick Hill <[email protected]> * [Core][V1][TPU] Enable structured decoding on TPU V1 (vllm-project#16499) Signed-off-by: Chenyaaang <[email protected]> * [Bugfix] validate urls object for multimodal content parts (vllm-project#16990) Signed-off-by: Guillaume Calmettes <[email protected]> * add Dockerfile build vllm against torch nightly (vllm-project#16936) Signed-off-by: Yang Wang <[email protected]> * [Kernel][ROCM] Upstream prefix prefill speed up for vLLM V1 (vllm-project#13305) Signed-off-by: Sage Moore <[email protected]> Signed-off-by: root <[email protected]> Signed-off-by: Aleksandr Malyshev <[email protected]> Signed-off-by: root <[email protected]> Signed-off-by: maleksan85 <[email protected]> Signed-off-by: <> Co-authored-by: Sage Moore <[email protected]> Co-authored-by: root <[email protected]> Co-authored-by: Aleksandr Malyshev <[email protected]> Co-authored-by: qli88 <[email protected]> Co-authored-by: root <[email protected]> * [V1][DP] More robust DP/EP dummy request coordination (vllm-project#16277) Signed-off-by: Nick Hill <[email protected]> * [BugFix] Revert ROCm Custom Paged Attention Env Flag Check (vllm-project#17022) Signed-off-by: vllmellm <[email protected]> * Revert "[Misc] Add S3 environment variables for better support of MinIO." (vllm-project#17021) * [misc] tune some env vars for GB200 (vllm-project#16992) Signed-off-by: youkaichao <[email protected]> * [INTEL-HPU][v0] Port delayed sampling to upstream (vllm-project#16949) Signed-off-by: Michal Adamczyk <[email protected]> Signed-off-by: Chendi Xue <[email protected]> Co-authored-by: Michal Adamczyk <[email protected]> * [doc] add download path tips (vllm-project#17013) Signed-off-by: reidliu41 <[email protected]> Co-authored-by: reidliu41 <[email protected]> * [Bugfix] Triton FA function takes no keyword arguments (vllm-project#16902) Signed-off-by: vllmellm <[email protected]> * [V1] Avoid socket errors during shutdown when requests are in in-flight (vllm-project#16807) Signed-off-by: Nick Hill <[email protected]> * [BugFix] llama4 fa3 fix - RuntimeError: scheduler_metadata must have shape (metadata_size) (vllm-project#16998) Signed-off-by: Lucas Wilkinson <[email protected]> * [Misc] Improve readability of get_open_port function. (vllm-project#17024) Signed-off-by: gitover22 <[email protected]> * [Bugfix] Fix AssertionError: skip_special_tokens=False is not supported for Mistral tokenizers (vllm-project#16964) Signed-off-by: chaunceyjiang <[email protected]> * [CI] Run v1/test_serial_utils.py in CI (vllm-project#16996) Signed-off-by: Russell Bryant <[email protected]> * Mistral-format support for compressed-tensors (vllm-project#16803) Signed-off-by: mgoin <[email protected]> * Categorize `tests/kernels/` based on kernel type (vllm-project#16799) Signed-off-by: mgoin <[email protected]> * [Doc] Add top anchor and a note to quantization/bitblas.md (vllm-project#17042) Signed-off-by: windsonsea <[email protected]> * Ensure that `pid` passed to `kill_process_tree` is `int` for `mypy` (vllm-project#17051) Signed-off-by: Harry Mellor <[email protected]> * [CI] Update structured-output label automation (vllm-project#17055) Signed-off-by: Russell Bryant <[email protected]> * Improve Transformers backend model loading QoL (vllm-project#17039) Signed-off-by: Harry Mellor <[email protected]> * `CacheConfig.block_size` should always be `int` when used (vllm-project#17052) Signed-off-by: Harry Mellor <[email protected]> * Use `@property` and private field for `data_parallel_rank_local` (vllm-project#17053) Signed-off-by: Harry Mellor <[email protected]> * [Frontend] Support guidance:no-additional-properties for compatibility with xgrammar (vllm-project#15949) Signed-off-by: Travis Johnson <[email protected]> * [BugFix][V1] Fix int32 token index overflow when preparing input ids (vllm-project#16806) * [V1][Spec Decode] Always use argmax for sampling draft tokens (vllm-project#16899) Signed-off-by: Woosuk Kwon <[email protected]> * [CI/Build] workaround for CI build failure (vllm-project#17070) Signed-off-by: csy1204 <[email protected]> Co-authored-by: Michael Goin <[email protected]> * [Quantization]add prefix for commandA quantized model (vllm-project#17017) * [Minor] Use larger batch sizes for A100/B100/B200/MI300x (vllm-project#17073) Signed-off-by: Woosuk Kwon <[email protected]> * [Bugfix] Enable V1 usage stats (vllm-project#16986) Signed-off-by: mgoin <[email protected]> Signed-off-by: Nick Hill <[email protected]> Co-authored-by: Nick Hill <[email protected]> * More informative error when using Transformers backend (vllm-project#16988) Signed-off-by: Harry Mellor <[email protected]> * Addendum Fix to support FIPS enabled machines with MD5 hashing (vllm-project#17043) Signed-off-by: sydarb <[email protected]> * [Bugfix][Core] add seq_id_to_seq_group clearing to avoid memory leak when s… (vllm-project#16472) Signed-off-by: 开哲 <[email protected]> Co-authored-by: 开哲 <[email protected]> * [V1] Update structured output (vllm-project#16812) Signed-off-by: reidliu41 <[email protected]> Co-authored-by: reidliu41 <[email protected]> * [doc] update to hyperlink (vllm-project#17096) Signed-off-by: reidliu41 <[email protected]> Co-authored-by: reidliu41 <[email protected]> * Add docs for runai_streamer_sharded (vllm-project#17093) Signed-off-by: Omer Dayan (SW-GPU) <[email protected]> Co-authored-by: Cyrus Leung <[email protected]> * [Chore] Remove Sampler from Model Code (vllm-project#17084) Signed-off-by: Woosuk Kwon <[email protected]> * Disable enforce_eager for V1 TPU sampler and structured output tests (vllm-project#17016) Signed-off-by: mgoin <[email protected]> * Simplify `TokenizerGroup` (vllm-project#16790) Signed-off-by: Harry Mellor <[email protected]> * Fix OOT registration test (vllm-project#17099) Signed-off-by: Harry Mellor <[email protected]> * [V1][PP] Optimization: continue scheduling prefill chunks (vllm-project#17080) Signed-off-by: Rui Qiao <[email protected]> * [Misc] Remove OLMo2 config copy (vllm-project#17066) Signed-off-by: Isotr0py <[email protected]> * Improve static type checking in `LoRAModelRunnerMixin` (vllm-project#17104) Signed-off-by: Harry Mellor <[email protected]> * [V1][Structured Output] Clear xgrammar compiler object when engine core shut down to avoid nanobind leaked warning (vllm-project#16954) Signed-off-by: shen-shanshan <[email protected]> * [Frontend] Using matryoshka_dimensions control the allowed output dimensions. (vllm-project#16970) * Add missing rocm_skinny_gemms kernel test to CI (vllm-project#17060) Signed-off-by: mgoin <[email protected]> * [Misc] refactor example series - structured outputs (vllm-project#17040) Signed-off-by: reidliu41 <[email protected]> Co-authored-by: reidliu41 <[email protected]> * [V1][Spec Decoding] Add num_drafts and num_accepted_tokens_per_position metrics (vllm-project#16665) Signed-off-by: Mark McLoughlin <[email protected]> * [CI] Add automation for the `tool-calling` github label (vllm-project#17118) Signed-off-by: Russell Bryant <[email protected]> * Updating builkite job for IBM Power (vllm-project#17111) Signed-off-by: Aaruni Aggarwal <[email protected]> * existing torch installation pip command fix for docs (vllm-project#17059) * Molmo Requirements (vllm-project#17026) Signed-off-by: Eyshika Agarwal <[email protected]> Signed-off-by: eyshika <[email protected]> * Add `:markdownhelp:` to `EngineArgs` docs so markdown docstrings render properly (vllm-project#17124) Signed-off-by: Harry Mellor <[email protected]> * Improve configs - `LoRAConfig` + `PromptAdapterConfig` (vllm-project#16980) Signed-off-by: Harry Mellor <[email protected]> * [Docs] Generate correct github links for decorated functions (vllm-project#17125) Signed-off-by: Russell Bryant <[email protected]> * Add collective_rpc to llm engine (vllm-project#16999) Signed-off-by: Yinghai Lu <[email protected]> * Add chat template for Llama 4 models (vllm-project#16428) Signed-off-by: Max de Bayser <[email protected]> * [Misc] Add example to run DeepSeek with Ray Serve LLM (vllm-project#17134) Signed-off-by: Rui Qiao <[email protected]> * Better error message for missing mistral params.json (vllm-project#17132) Signed-off-by: mgoin <[email protected]> * Use custom address for listening socket (vllm-project#15988) Signed-off-by: Jens Glaser <[email protected]> * [FEAT] [ROCm]: AITER Fused MOE V1 Support (vllm-project#16752) Signed-off-by: vllmellm <[email protected]> Co-authored-by: tjtanaa <[email protected]> * [Attention] FA3 decode perf improvement - single mma warp group support for head dim 128 (vllm-project#16864) Signed-off-by: Lucas Wilkinson <[email protected]> * fix float16 support for kimi-vl (vllm-project#17156) Co-authored-by: zhouzaida <[email protected]> * [Doc] V1 : Update LoRA status (vllm-project#17133) Signed-off-by: varun sundar rabindranath <[email protected]> Co-authored-by: varun sundar rabindranath <[email protected]> * [Docs] Fix True->true in supported_models.md (vllm-project#17141) * Move missed `SchedulerConfig` args into scheduler config group in `EngineArgs` (vllm-project#17131) Signed-off-by: Harry Mellor <[email protected]> * [Misc] Clean up redundant code in uniproc_executor.py (vllm-project#16762) Signed-off-by: Lifu Huang <[email protected]> * [Bugfix][Misc] Use TritonPlaceholderModule to defensively import triton (vllm-project#15099) Signed-off-by: Mengqing Cao <[email protected]> * [Misc] Benchmark Serving Script Support Appending Results (vllm-project#17028) Signed-off-by: Lucas Wilkinson <[email protected]> * [Perf]Optimize rotary_emb implementation to use Triton operator for improved inference performance (vllm-project#16457) Signed-off-by: cynthieye <[email protected]> Co-authored-by: MagnetoWang <[email protected]> * [Bugfix] remove fallback in guided_json (int range, patterns) (vllm-project#16725) Signed-off-by: csy1204 <[email protected]> Co-authored-by: 조상연[플레이스 AI] <[email protected]> * [Quantization][FP8] Add support for FP8 models with input_scale for output projection and QK quantization (vllm-project#15734) Signed-off-by: Randall Smith <[email protected]> Signed-off-by: Luka Govedič <[email protected]> Co-authored-by: Luka Govedič <[email protected]> * [Doc] Add headings to improve gptqmodel.md (vllm-project#17164) Signed-off-by: windsonsea <[email protected]> * Only turn on FastIncrementalDetokenizer when tokenizers >= 0.21.1 (vllm-project#17158) * [Doc] Add two links to disagg_prefill.md (vllm-project#17168) Signed-off-by: windsonsea <[email protected]> * [Doc] Move todo out of beam search docstring (vllm-project#17183) Signed-off-by: Alex-Brooks <[email protected]> * [Bugfix] Fix mistral model tests (vllm-project#17181) Signed-off-by: DarkLight1337 <[email protected]> * [Bugfix] Fix Mistral ChatCompletionRequest Body Exception (vllm-project#16769) Signed-off-by: Jasmond Loh <[email protected]> Co-authored-by: Cyrus Leung <[email protected]> * Fix API typo and remove FP8 on V1 restriction --------- Signed-off-by: Nick Hill <[email protected]> Signed-off-by: Chenyaaang <[email protected]> Signed-off-by: Guillaume Calmettes <[email protected]> Signed-off-by: Yang Wang <[email protected]> Signed-off-by: Sage Moore <[email protected]> Signed-off-by: root <[email protected]> Signed-off-by: Aleksandr Malyshev <[email protected]> Signed-off-by: root <[email protected]> Signed-off-by: maleksan85 <[email protected]> Signed-off-by: <> Signed-off-by: vllmellm <[email protected]> Signed-off-by: youkaichao <[email protected]> Signed-off-by: Michal Adamczyk <[email protected]> Signed-off-by: Chendi Xue <[email protected]> Signed-off-by: reidliu41 <[email protected]> Signed-off-by: Lucas Wilkinson <[email protected]> Signed-off-by: gitover22 <[email protected]> Signed-off-by: chaunceyjiang <[email protected]> Signed-off-by: Russell Bryant <[email protected]> Signed-off-by: mgoin <[email protected]> Signed-off-by: windsonsea <[email protected]> Signed-off-by: Harry Mellor <[email protected]> Signed-off-by: Travis Johnson <[email protected]> Signed-off-by: Woosuk Kwon <[email protected]> Signed-off-by: csy1204 <[email protected]> Signed-off-by: sydarb <[email protected]> Signed-off-by: 开哲 <[email protected]> Signed-off-by: Omer Dayan (SW-GPU) <[email protected]> Signed-off-by: Rui Qiao <[email protected]> Signed-off-by: Isotr0py <[email protected]> Signed-off-by: shen-shanshan <[email protected]> Signed-off-by: Mark McLoughlin <[email protected]> Signed-off-by: Aaruni Aggarwal <[email protected]> Signed-off-by: Eyshika Agarwal <[email protected]> Signed-off-by: eyshika <[email protected]> Signed-off-by: Yinghai Lu <[email protected]> Signed-off-by: Max de Bayser <[email protected]> Signed-off-by: Jens Glaser <[email protected]> Signed-off-by: varun sundar rabindranath <[email protected]> Signed-off-by: Lifu Huang <[email protected]> Signed-off-by: Mengqing Cao <[email protected]> Signed-off-by: cynthieye <[email protected]> Signed-off-by: Randall Smith <[email protected]> Signed-off-by: Luka Govedič <[email protected]> Signed-off-by: Alex-Brooks <[email protected]> Signed-off-by: DarkLight1337 <[email protected]> Signed-off-by: Jasmond Loh <[email protected]> Co-authored-by: Nick Hill <[email protected]> Co-authored-by: Chenyaaang <[email protected]> Co-authored-by: Guillaume Calmettes <[email protected]> Co-authored-by: Yang Wang <[email protected]> Co-authored-by: Aleksandr Malyshev <[email protected]> Co-authored-by: Sage Moore <[email protected]> Co-authored-by: root <[email protected]> Co-authored-by: Aleksandr Malyshev <[email protected]> Co-authored-by: qli88 <[email protected]> Co-authored-by: root <[email protected]> Co-authored-by: vllmellm <[email protected]> Co-authored-by: Chauncey <[email protected]> Co-authored-by: youkaichao <[email protected]> Co-authored-by: Chendi.Xue <[email protected]> Co-authored-by: Michal Adamczyk <[email protected]> Co-authored-by: Reid <[email protected]> Co-authored-by: reidliu41 <[email protected]> Co-authored-by: Lucas Wilkinson <[email protected]> Co-authored-by: huafeng <[email protected]> Co-authored-by: Russell Bryant <[email protected]> Co-authored-by: Michael Goin <[email protected]> Co-authored-by: Michael Yao <[email protected]> Co-authored-by: Harry Mellor <[email protected]> Co-authored-by: Travis Johnson <[email protected]> Co-authored-by: Yong Hoon Shin <[email protected]> Co-authored-by: Woosuk Kwon <[email protected]> Co-authored-by: Sangyeon Cho <[email protected]> Co-authored-by: Chen Xia <[email protected]> Co-authored-by: Areeb Syed <[email protected]> Co-authored-by: 张宇 <[email protected]> Co-authored-by: 开哲 <[email protected]> Co-authored-by: omer-dayan <[email protected]> Co-authored-by: Cyrus Leung <[email protected]> Co-authored-by: Rui Qiao <[email protected]> Co-authored-by: Isotr0py <[email protected]> Co-authored-by: Shanshan Shen <[email protected]> Co-authored-by: wang.yuqi <[email protected]> Co-authored-by: Mark McLoughlin <[email protected]> Co-authored-by: Aaruni Aggarwal <[email protected]> Co-authored-by: Atilla <[email protected]> Co-authored-by: Eyshika Agarwal <[email protected]> Co-authored-by: Yinghai Lu <[email protected]> Co-authored-by: Maximilien de Bayser <[email protected]> Co-authored-by: jglaser <[email protected]> Co-authored-by: tjtanaa <[email protected]> Co-authored-by: Zaida Zhou <[email protected]> Co-authored-by: zhouzaida <[email protected]> Co-authored-by: Varun Sundar Rabindranath <[email protected]> Co-authored-by: varun sundar rabindranath <[email protected]> Co-authored-by: Lifu Huang <[email protected]> Co-authored-by: Mengqing Cao <[email protected]> Co-authored-by: yexin(叶鑫) <[email protected]> Co-authored-by: MagnetoWang <[email protected]> Co-authored-by: 조상연[플레이스 AI] <[email protected]> Co-authored-by: rasmith <[email protected]> Co-authored-by: Luka Govedič <[email protected]> Co-authored-by: Lu Fang <[email protected]> Co-authored-by: Alex Brooks <[email protected]> Co-authored-by: Cyrus Leung <[email protected]> Co-authored-by: Jasmond L <[email protected]>
…mproved inference performance (vllm-project#16457) Signed-off-by: cynthieye <[email protected]> Co-authored-by: MagnetoWang <[email protected]> Signed-off-by: Zijing Liu <[email protected]>
…mproved inference performance (vllm-project#16457) Signed-off-by: cynthieye <[email protected]> Co-authored-by: MagnetoWang <[email protected]>
…mproved inference performance (vllm-project#16457) Signed-off-by: cynthieye <[email protected]> Co-authored-by: MagnetoWang <[email protected]>
…mproved inference performance (vllm-project#16457) Signed-off-by: cynthieye <[email protected]> Co-authored-by: MagnetoWang <[email protected]>
Optimize Rotary Positional Embeddings with Triton Kernel in VLLM
This PR enhances rotary positional embedding computation by leveraging Triton-optimized kernels from flash_attn, addressing a significant performance bottleneck observed in models like Qwen2-VL.
Background
The original PyTorch-native rotary embedding implementation (rotary_emb) consumed 40-60% of total inference latency for Qwen2-VL, particularly scaling with output token count. Profiling revealed inefficiencies in tensor reshaping and element-wise operations.
Changes
Triton Kernel Integration:
Conditionally uses flash_attn.ops.triton.rotary.apply_rotary when flash-attn>=2.0 is available.
Falls back to the native PyTorch implementation otherwise.
Dynamic Implementation Selection:
Added _use_flash_attn flag to RotaryEmbedding classes, auto-detecting flash_attn availability during initialization.
Performance Gains
Qwen2-VL: Achieved 17% end-to-end speedup when generating 150 output tokens.
Generalization: Expected improvements for other mrotary-based models, pending further benchmarks.
Code Compatibility
Maintains backward compatibility with existing model architectures.
Requires no user-side changes—automatically prioritizes Triton kernel when dependencies are met.
This optimization significantly reduces rotary embedding overhead while preserving numerical equivalence, making it particularly impactful for long-context and long-generation scenarios.