-
-
Notifications
You must be signed in to change notification settings - Fork 10.5k
Add CUDA graph-based all reduce launcher #26
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
Conversation
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM!
self.group = get_tensor_model_parallel_group() | ||
self.buffer = torch.empty( | ||
size=(max_num_tokens, hidden_size), | ||
dtype=torch.half, # FIXME: hardcoded dtype |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Add a dtype argument for this class?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Fixed!
Disable NPU merged to OV master recently
Install and configure use of the NCCL version recommended by vLLM via the [vllm-nccl](https://github.com/vllm-project/vllm-nccl) package. The install is a little wonky... but this set of changes should work. Signed-off-by: Travis Johnson <[email protected]>
deps: bump fastapi to >= 0.109.1
Update max_context_len for custom paged attention.
…c466a3 Rebase habana_main up to cc466a3
…inear_fusion_and_prepack Enable linear fusion/prepack and MOE AWQ fusion
* add tool server Signed-off-by: Chen Zhang <[email protected]> * add back demo tool server Signed-off-by: Chen Zhang <[email protected]> * update Signed-off-by: Chen Zhang <[email protected]> * update Signed-off-by: Chen Zhang <[email protected]> * update disallow cases Signed-off-by: Chen Zhang <[email protected]> * fix Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> --------- Signed-off-by: Chen Zhang <[email protected]>
* add tool server Signed-off-by: Chen Zhang <[email protected]> * add back demo tool server Signed-off-by: Chen Zhang <[email protected]> * update Signed-off-by: Chen Zhang <[email protected]> * update Signed-off-by: Chen Zhang <[email protected]> * update disallow cases Signed-off-by: Chen Zhang <[email protected]> * fix Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> * fix some type Signed-off-by: Chen Zhang <[email protected]> --------- Signed-off-by: Chen Zhang <[email protected]>
…oject#26) * indexer medatata to separate prefill and decode * deep_gemm prefill kernel * decode kernel, can run for single batch * bug fixing insert decode k into kv before gemm * don't use tilelang quant function * faster non-looping torch for kv cache insertion * add chunked prefill impl * change quant kernel back to tilelang for promotion * fix format (vllm-project#31) Signed-off-by: Chen Zhang <[email protected]> * update unit tests * Fp8 indexer prefill (vllm-project#33) * init Signed-off-by: Chen Zhang <[email protected]> * can run --------- Signed-off-by: Chen Zhang <[email protected]> * remove debug comment Signed-off-by: Chen Zhang <[email protected]> * cleanup * further cleanup --------- Signed-off-by: Chen Zhang <[email protected]> Co-authored-by: mgoin <[email protected]> Co-authored-by: Chen Zhang <[email protected]>
Related to #22
This PR uses CUDA graph to reduce the CPU overhead of NCCL all reduce operation.