Skip size calculation during async copy wait=True #1088
Closed
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Signed-off-by: Rupeng Liu rupengliu@meta.com
Contributor
|
around 5-10% throughput improvement |
bythew3i
reviewed
Nov 19, 2025
| sem, | ||
| wait, | ||
| src=vmem_ref, | ||
| dst=vmem_ref.at[pl.ds(0, offset + bkv_sz_frm_new)], |
Collaborator
There was a problem hiding this comment.
wait only cares about size.
Please refer to https://github.com/vllm-project/tpu-inference/pull/718/files
Contributor
|
moved to a new github account, close this PR for now |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Thanks to @bythew3i (Jevin) for providing the insight for the optimization. This is to skip size computation during DMA wait.
I have ran both kernel's test and brought up e2e vllm server to test, will provide perf improvement numbers
Signed-off-by: Rupeng Liu rupengliu@meta.com
Description
Start with a short description of what the PR does and how this is a change from
the past.
The rest of the description includes relevant details and context, examples:
If the change fixes a bug or a Github issue, please include a link, e.g.,:
FIXES: b/123456
FIXES: #123456
Tests
Please describe how you tested this change, and include any instructions and/or
commands to reproduce.
Checklist
Before submitting this PR, please make sure: