Skip to content

[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

Merged
merged 2 commits into from
Apr 25, 2025

Conversation

cynthieye
Copy link
Contributor

@cynthieye cynthieye commented Apr 11, 2025

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.

Copy link

👋 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 fastcheck CI which starts running only a small and essential subset of CI tests to quickly catch errors. You can run other CI tests on top of those by going to your fastcheck build on Buildkite UI (linked in the PR checks section) and unblock them. If you do not have permission to unblock, ping simon-mo or khluu to add you in our Buildkite org.

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 ready label to the PR or enable auto-merge.

🚀

@cynthieye cynthieye force-pushed the main branch 2 times, most recently from a115f1f to e630289 Compare April 11, 2025 06:59
@jeejeelee
Copy link
Collaborator

Perhaps using kernel from https://github.com/vllm-project/flash-attention would be more reasonable

@MagnetoWang MagnetoWang force-pushed the main branch 2 times, most recently from 39b1672 to b2379fd Compare April 16, 2025 07:34
@cynthieye cynthieye force-pushed the main branch 6 times, most recently from c435afc to 8d6303f Compare April 18, 2025 06:34
@Isotr0py
Copy link
Collaborator

Shouldn't this triton kernel have been used for Qwen2-VL?

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

use_flash_attn = self.attn_backend == _Backend.FLASH_ATTN
q = apply_rotary_pos_emb_vision(q,
rotary_pos_emb,
use_flash_attn=use_flash_attn)
k = apply_rotary_pos_emb_vision(k,
rotary_pos_emb,
use_flash_attn=use_flash_attn)

@cynthieye
Copy link
Contributor Author

cynthieye commented Apr 18, 2025

apply_rotary_pos_emb_vision

qwen2_vl.py

class Qwen2VLForConditionalGeneration(nn.Module, SupportsMultiModal,
                                      SupportsLoRA, SupportsPP):
    def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""):
        super().__init__()
        ......

        self.language_model = init_vllm_registered_model(
            vllm_config=vllm_config,
            prefix=maybe_prefix(prefix, "language_model"),
            architectures=["Qwen2ForCausalLM"],
        )
        ......
       
            
    def forward(
        self,
        input_ids: torch.Tensor,
        positions: torch.Tensor,
        intermediate_tensors: Optional[IntermediateTensors] = None,
        inputs_embeds: Optional[torch.Tensor] = None,
        **kwargs: object,
    ) -> Union[torch.Tensor, IntermediateTensors]:
        ......

        hidden_states = self.language_model.model(
            input_ids=input_ids,
            positions=positions,
            intermediate_tensors=intermediate_tensors,
            inputs_embeds=inputs_embeds,
        )
        return hidden_states

qwen2.py

from vllm.model_executor.layers.rotary_embedding import get_rope

class Qwen2Attention(nn.Module):
    def __init__(self,
                 hidden_size: int,
                 num_heads: int,
                 num_kv_heads: int,
                 max_position: int = 4096 * 32,
                 rope_theta: float = 10000,
                 cache_config: Optional[CacheConfig] = None,
                 quant_config: Optional[QuantizationConfig] = None,
                 rope_scaling: Optional[Tuple] = None) -> None:
        super().__init__()
        ......

        self.rotary_emb = get_rope(
            self.head_dim,
            rotary_dim=self.head_dim,
            max_position=max_position,
            base=self.rope_theta,
            rope_scaling=rope_scaling,
        )
        ......
    def forward(
        self,
        positions: torch.Tensor,
        hidden_states: torch.Tensor,
        kv_cache: torch.Tensor,
        attn_metadata: AttentionMetadata,
    ) -> torch.Tensor:
        ......
        q, k = self.rotary_emb(positions, q, k)
        ......

@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
Uploading screenshot-20250418-203723.png…

Comment on lines 88 to 85
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)
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
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.

Copy link
Contributor Author

@cynthieye cynthieye Apr 19, 2025

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

Copy link
Contributor Author

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

Copy link
Collaborator

@Isotr0py Isotr0py Apr 19, 2025

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.

Copy link
Contributor Author

@cynthieye cynthieye Apr 19, 2025

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'

Copy link
Contributor Author

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:

  1. Create a 'layers' sub directory in the' vllm_flash-attn 'directory
  2. Then copy 'flash_attn/layers/total. py' to the 'vllm_flash-attn/layers/' directory,
  3. Import or export 'apply_rotary-emb' in 'vllm_flash-attn/layers/init. py'.
  4. Ensure that the top-level 'init. py' of 'vllm_flash-attn' correctly exports the 'layers' module.

Copy link
Collaborator

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.

Copy link
Contributor Author

@cynthieye cynthieye Apr 20, 2025

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

Copy link
Contributor Author

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

Copy link
Contributor Author

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

@cynthieye cynthieye force-pushed the main branch 4 times, most recently from 4dca833 to 679db29 Compare April 18, 2025 17:39
@cynthieye cynthieye requested a review from Isotr0py April 18, 2025 21:16
Isotr0py
Isotr0py previously approved these changes Apr 19, 2025
Copy link
Collaborator

@Isotr0py Isotr0py left a 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!

@Isotr0py Isotr0py added the ready ONLY add when PR is ready to merge/full CI is needed label Apr 19, 2025
@cynthieye cynthieye force-pushed the main branch 4 times, most recently from 73354b5 to bf7c6fc Compare April 19, 2025 18:07
Copy link

mergify bot commented Apr 25, 2025

This pull request has merge conflicts that must be resolved before it can be
merged. Please rebase the PR, @cynthieye.

https://docs.github.com/en/pull-requests/collaborating-with-pull-requests/working-with-forks/syncing-a-fork

@mergify mergify bot added the needs-rebase label Apr 25, 2025
@cynthieye cynthieye force-pushed the main branch 2 times, most recently from 0069947 to c04e840 Compare April 25, 2025 04:05
@mergify mergify bot removed the needs-rebase label Apr 25, 2025
…mproved inference performance

Signed-off-by: cynthieye <[email protected]>
Co-authored-by: MagnetoWang <[email protected]>
Copy link

mergify bot commented Apr 25, 2025

This pull request has merge conflicts that must be resolved before it can be
merged. Please rebase the PR, @cynthieye.

https://docs.github.com/en/pull-requests/collaborating-with-pull-requests/working-with-forks/syncing-a-fork

…mproved inference performance

Signed-off-by: cynthieye <[email protected]>
Co-authored-by: MagnetoWang <[email protected]>
Copy link
Collaborator

@Isotr0py Isotr0py left a 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!

@Isotr0py Isotr0py merged commit b22980a into vllm-project:main Apr 25, 2025
44 checks passed
gshtras added a commit to ROCm/vllm that referenced this pull request Apr 25, 2025
* [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]>
liuzijing2014 pushed a commit to liuzijing2014/vllm that referenced this pull request Apr 25, 2025
…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]>
liuzijing2014 pushed a commit to liuzijing2014/vllm that referenced this pull request Apr 25, 2025
…mproved inference performance (vllm-project#16457)

Signed-off-by: cynthieye <[email protected]>
Co-authored-by: MagnetoWang <[email protected]>
wuisawesome pushed a commit to character-tech/vllm that referenced this pull request Apr 28, 2025
…mproved inference performance (vllm-project#16457)

Signed-off-by: cynthieye <[email protected]>
Co-authored-by: MagnetoWang <[email protected]>
jikunshang pushed a commit to jikunshang/vllm that referenced this pull request Apr 29, 2025
…mproved inference performance (vllm-project#16457)

Signed-off-by: cynthieye <[email protected]>
Co-authored-by: MagnetoWang <[email protected]>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
ci/build ready ONLY add when PR is ready to merge/full CI is needed
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants