[Bugfix][Kernel] allow non-power-of-2 for prefix prefill with alibi #4573
Conversation
|
LGTM FWIW! |
ok, I will add a test for prefix prefill kernel with alibi later. |
|
The local test passed. Waiting for all kernel tests of CI to pass. pytest test_prefix_prefill.py
================================================================================== test session starts ===================================================================================
platform linux -- Python 3.10.13, pytest-8.1.0, pluggy-1.4.0
configfile: pyproject.toml
plugins: shard-0.1.2, rerunfailures-13.0, anyio-4.3.0, asyncio-0.23.5, forked-1.6.0
asyncio: mode=strict
collected 48 items
Running 48 items in this shard
test_prefix_prefill.py ................................................ [100%]
============================================================================= 48 passed in 157.21s (0:02:37) ============================================================================= |
|
sorry I am going to take a look at it today! |
|
@rkooo567 alibi tests have been moved to test_prefix_prefill.py. Kernel tests in CI have passed. |
There was a problem hiding this comment.
So the logic change itself is equivalent to non-alibi case right?
Also, @WoosukKwon it'd be great if you can take a look at alibi slope testing. It looks okay to me, but just in case...
yes, the logic is the same as non-alibi case. The additional dim_mask collaborate with offs_m can generate right mask for tl.load. for example: # dim_mask[None, :] & offs_m[:, None] -> [1,D] & [M, 1]
# >>> dim_mask = torch.tensor([1,1,1,1,0,0]) # [1,6]
# >>> offs_m = torch.tensor([8,9,10,11]) # [4,1]
# >>> mask=dim_mask[None, :] & (offs_m[:, None] < 12)
# >>> mask
# tensor([[1, 1, 1, 1, 0, 0],
# [1, 1, 1, 1, 0, 0],
# [1, 1, 1, 1, 0, 0],
# [1, 1, 1, 1, 0, 0]])
# >>> mask.shape
# torch.Size([4, 6]) # [M,D]
q = tl.load(Q + off_q, # [M,D]
mask=dim_mask[None, :] &
(offs_m[:, None] < cur_batch_seq_len - cur_batch_ctx_len),
other=0.0)I can add some comments if needed. |
Head branch was pushed to by a user without write access
|
@rkooo567 seems many tests interrupted by a signal |
|
can you try merge the latest master? I saw it sometimes happens... not sure what's the root cause |
FILL IN THE PR DESCRIPTION HERE
FIX #4171
allow non-power-of-two head sizes in prefix prefill with alibi, this is a small fix based on #4128.
BEFORE SUBMITTING, PLEASE READ THE CHECKLIST BELOW AND FILL IN THE DESCRIPTION ABOVE
PR Checklist (Click to Expand)
Thank you for your contribution to vLLM! Before submitting the pull request, please ensure the PR meets the following criteria. This helps vLLM maintain the code quality and improve the efficiency of the review process.
PR Title and Classification
Only specific types of PRs will be reviewed. The PR title is prefixed appropriately to indicate the type of change. Please use one of the following:
[Bugfix]for bug fixes.[CI/Build]for build or continuous integration improvements.[Doc]for documentation fixes and improvements.[Model]for adding a new model or improving an existing model. Model name should appear in the title.[Frontend]For changes on the vLLM frontend (e.g., OpenAI API server,LLMclass, etc.)[Kernel]for changes affecting CUDA kernels or other compute kernels.[Core]for changes in the core vLLM logic (e.g.,LLMEngine,AsyncLLMEngine,Scheduler, etc.)[Hardware][Vendor]for hardware-specific changes. Vendor name should appear in the prefix (e.g.,[Hardware][AMD]).[Misc]for PRs that do not fit the above categories. Please use this sparingly.Note: If the PR spans more than one category, please include all relevant prefixes.
Code Quality
The PR need to meet the following code quality standards:
format.shto format your code.docs/source/if the PR modifies the user-facing behaviors of vLLM. It helps vLLM user understand and utilize the new features or changes.Notes for Large Changes
Please keep the changes as concise as possible. For major architectural changes (>500 LOC excluding kernel/data/config/test), we would expect a GitHub issue (RFC) discussing the technical design and justification. Otherwise, we will tag it with
rfc-requiredand might not go through the PR.What to Expect for the Reviews
The goal of the vLLM team is to be a transparent reviewing machine. We would like to make the review process transparent and efficient and make sure no contributor feel confused or frustrated. However, the vLLM team is small, so we need to prioritize some PRs over others. Here is what you can expect from the review process:
action-requiredlabel on the PR if there are changes required. The contributor should address the comments and ping the reviewer to re-review the PR.Thank You
Finally, thank you for taking the time to read these guidelines and for your interest in contributing to vLLM. Your contributions make vLLM a great tool for everyone!