Skip to content
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.

Sign up for GitHub

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

Jump to bottom

[Kernel] LoRA - Refactor sgmv kernels #13110

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
Open

[Kernel] LoRA - Refactor sgmv kernels #13110

wants to merge 1 commit into from
+327 −129

Conversation

varun-sundar-rabindranath
Copy link
Contributor

Refactor / Modularize sgmv_shrink and sgmv_expand kernels to enable better code re-use.

refactor sgmv kernels
024b2af 
Signed-off-by: Varun Sundar Rabindranath <[email protected]>
@github-actions GitHub Actions
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.

🚀

varun-sundar-rabindranath
varun-sundar-rabindranath commented Feb 11, 2025
View reviewed changes
vllm/lora/ops/triton_ops/sgmv_shrink.py
@@ -62,67 +63,50 @@ def _sgmv_shrink_kernel(
pid_sk = pid_mix % SPLIT_K

M = tl.load(seq_lens + cur_batch)
if pid_m * BLOCK_M > M:
if pid_m * BLOCK_M >= M:
Copy link
Contributor Author

Choose a reason for hiding this comment

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

calling attention to a functional change - I believe it was a bug before.

varun-sundar-rabindranath
varun-sundar-rabindranath commented Feb 11, 2025
View reviewed changes
vllm/lora/ops/triton_ops/sgmv_expand.py
return
if pid_n * BLOCK_N > curr_N:
if pid_n * BLOCK_N >= curr_N:
return
Copy link
Contributor Author

Choose a reason for hiding this comment

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

Calling attention to functional changes, I believe,
if pid_m * BLOCK_M > M: and
if pid_n * BLOCK_N > curr_N: were bugs before.

@varun-sundar-rabindranath
Copy link
Contributor Author

varun-sundar-rabindranath commented Feb 11, 2025 •
edited
Loading

I have a PR #13096 for V1 that introduces a shrink and expand kernels. The V1 kernels are very similar to the SGMV kernels. This refactor helps that PR also.
I have marked #13096 as draft for now - I'll open that PR once this PR lands 🙏

@jeejeelee PTAL ! Thanks !

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers
No reviews
Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
1 participant
@varun-sundar-rabindranath