[Kernel] LoRA - Refactor sgmv kernels #13110
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 🚀 |
| @@ -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: | |||
There was a problem hiding this comment.
calling attention to a functional change - I believe it was a bug before.
| return | ||
| if pid_n * BLOCK_N > curr_N: | ||
| if pid_n * BLOCK_N >= curr_N: | ||
| return |
There was a problem hiding this comment.
Calling attention to functional changes, I believe,
if pid_m * BLOCK_M > M: and
if pid_n * BLOCK_N > curr_N: were bugs before.
|
I have a PR #13096 for V1 that introduces a @jeejeelee PTAL ! Thanks ! |
Signed-off-by: Varun Sundar Rabindranath <[email protected]>
Signed-off-by: Varun Sundar Rabindranath <[email protected]>
varun-sundar-rabindranath
force-pushed
the
varun/sgmv-kernel-refactor
branch
from
3e94503 to
7d4b9fa
Compare
Signed-off-by: Michael Glass <[email protected]>
Signed-off-by: Linkun Chen <[email protected]>
Refactor / Modularize sgmv_shrink and sgmv_expand kernels to enable better code re-use.