-
Notifications
You must be signed in to change notification settings - Fork 432
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
profile throughput without new threads #2826
base: main
Are you sure you want to change the base?
Conversation
@@ -742,8 +688,9 @@ async def async_forward(self, inputs: ModelInputs, swap_in_map: SwapMap, | |||
output = self._forward_impl(inputs, | |||
swap_in_map=swap_in_map, | |||
swap_out_map=swap_out_map) | |||
await asyncio.get_event_loop().run_in_executor(None, | |||
self.stream.synchronize) | |||
await asyncio.sleep(0) |
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.
How does this method outperform the previous one?
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.
asyncio.sleep
release CPU without creating new thread
@@ -741,6 +748,8 @@ def __update_inputs(next_token_ids): | |||
logger.debug('<ForwardTask>: ' | |||
f'batch_size={inputs.seq_length.size(0)} ' | |||
f'num_tokens={inputs.input_ids.size(-1)}') | |||
if self.gpu_count == 1: |
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.
Why do we need this while previously not?
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.
logits process can be streaming with cuda inputs.
distribute container with cuda tensor is slower than CPU tensor.
@@ -60,8 +34,8 @@ def apply_rotary_pos_emb_qk_kernel( | |||
BLOCK_N: tl.constexpr, | |||
): | |||
"""apply rotary on key AND query kernel.""" | |||
seq_block_id = tl.program_id(0) | |||
head_id = tl.program_id(1) | |||
seq_block_id = tl.program_id(1) |
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.
What is the benefit of this change?
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.
increate cos/sin cache hit rate.
@@ -79,7 +47,7 @@ def _fill_kv_cache_kernel( | |||
QSeqLens, | |||
KVSeqLens, | |||
BlockOffsets, | |||
num_heads: tl.constexpr, | |||
is_decoding: tl.constexpr, |
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.
Does it mean that the quant kernel of fill_kv_cache
can also be optimized?
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.
quant kernel would be optimized in future.
Thanks for your contribution and we appreciate it a lot. The following instructions would make your pull request more healthy and more easily receiving feedbacks. If you do not understand some items, don't worry, just make the pull request and seek help from maintainers.
Motivation
Please describe the motivation of this PR and the goal you want to achieve through this PR.
Modification
Please briefly describe what modification is made in this PR.
BC-breaking (Optional)
Does the modification introduce changes that break the backward-compatibility of the downstream repositories?
If so, please describe how it breaks the compatibility and how the downstream projects should modify their code to keep compatibility with this PR.
Use cases (Optional)
If this PR introduces a new feature, it is better to list some use cases here, and update the documentation.
Checklist