[POC]Segmented spans#2
Conversation
Signed-off-by: David Ben-David <davidb@pliops.com>
…ed-token gaps Signed-off-by: David Ben-David <davidb@pliops.com>
Signed-off-by: Kfir Wolfson <kfirw@pliops.com>
Signed-off-by: Omer Paz <omerpaz95@gmail.com>
|
👋 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 You ask your reviewers to trigger select CI tests on top of 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 If you have any questions, please reach out to us on Slack at https://slack.vllm.ai. 🚀 |
Signed-off-by: omerpaz95 <omerpaz95@gmail.com>
Signed-off-by: omerpaz95 <omerpaz95@gmail.com>
Signed-off-by: omerpaz95 <omerpaz95@gmail.com>
There was a problem hiding this comment.
Pull request overview
This pull request introduces a proof-of-concept implementation for "Segmented Spans" functionality in vLLM, which enables handling of non-contiguous KV cache through virtual gap requests and span-based attention. The implementation integrates with the Legolink system and adds support for handling gaps in external KV cache by creating virtual requests for recomputation.
Changes:
- Added segmented prefill support through virtual gap request mechanism that allows recomputation of missing KV cache blocks
- Implemented fused RoPE (Rotary Position Embedding) in Triton attention kernels to improve performance
- Extended KV connector API with
get_computed_token_gaps()method to report non-contiguous cache hits - Added environment variables and configuration for span detection and block attention features
Reviewed changes
Copilot reviewed 30 out of 30 changed files in this pull request and generated 27 comments.
Show a summary per file
| File | Description |
|---|---|
| vllm/v1/core/sched/scheduler.py | Implements virtual gap request creation and scheduling logic for segmented prefill |
| vllm/v1/core/sched/output.py | Adds data structures for gap recomputation metadata and virtual request tracking |
| vllm/v1/core/sched/async_scheduler.py | Updates async scheduler to handle virtual gap requests |
| vllm/v1/worker/gpu_model_runner.py | Adds KV cache hash debugging, virtual request slot mapping, and cleanup logic |
| vllm/v1/worker/gpu/model_runner.py | Extends execute/sample workflow to track and cleanup virtual gap requests |
| vllm/v1/attention/ops/triton_unified_attention.py | Implements fused RoPE in attention kernels by splitting Q/K into halves |
| vllm/v1/attention/backends/triton_attn.py | Passes cos_sin_cache to attention ops when spans are enabled |
| vllm/v1/attention/backend.py | Adds cos_sin_cache field to CommonAttentionMetadata |
| vllm/model_executor/layers/rotary_embedding/base.py | Conditionally skips key RoPE application when spans are enabled |
| vllm/v1/core/kv_cache_utils.py | Implements span separator token detection for fan-in and recompute token handling |
| vllm/v1/core/block_pool.py | Adds deduplication logic when freeing blocks to handle shared blocks |
| vllm/envs.py | Adds environment variables for span configuration (VLLM_V1_SPANS_*) |
| vllm/distributed/kv_transfer/kv_connector/v1/base.py | Adds get_computed_token_gaps() API method to base connector |
| vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py | Adds null checks for virtual requests in offloading connector |
| vllm/distributed/kv_transfer/kv_connector/v1/example_connector.py | Fixes shape indexing bugs and adds req_id tracking to ReqMeta |
| examples/offline_inference/spans/* | Adds example scripts and connectors demonstrating spans functionality |
| examples/offline_inference/kv_segmented_prefill/* | Adds test framework for validating segmented prefill correctness |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| firstok = curr_block_token_ids[0] | ||
| if firstok == envs.VLLM_V1_SPANS_TOKEN_PLUS: |
There was a problem hiding this comment.
The variable name 'firstok' appears to be a typo for 'first_tok' or 'first_token'. This should be renamed for clarity.
| firstok = curr_block_token_ids[0] | |
| if firstok == envs.VLLM_V1_SPANS_TOKEN_PLUS: | |
| first_tok = curr_block_token_ids[0] | |
| if first_tok == envs.VLLM_V1_SPANS_TOKEN_PLUS: |
|
|
||
| extra_keys = recompute_token_handler( | ||
| block_tokens[0], block_tokens[:start_token_idx], extra_keys) | ||
|
|
There was a problem hiding this comment.
There is inconsistent trailing whitespace after this line. Trailing whitespace should be removed for code cleanliness.
| extra_keys = recompute_token_handler( | |
| block_tokens[0], block_tokens[:start_token_idx], extra_keys) | |
| extra_keys = recompute_token_handler( | |
| block_tokens[0], block_tokens[:start_token_idx], extra_keys | |
| ) |
| mask=dim_mask_b[None, :] & query_mask_0[:, None] & query_mask_1[:, None], | ||
| other=0.0, | ||
| ) | ||
|
|
There was a problem hiding this comment.
There is inconsistent trailing whitespace after this line. Trailing whitespace should be removed for code cleanliness.
| VLLM_LOG_MODEL_INSPECTION: bool = False | ||
| VLLM_DEBUG_MFU_METRICS: bool = False | ||
|
|
||
|
|
There was a problem hiding this comment.
There is inconsistent trailing whitespace after this line. Trailing whitespace should be removed for code cleanliness.
| print( | ||
| f"Choosing gaps. external_end = {external_end}, external_start = {external_start}" | ||
| ) | ||
| if external_end - external_start < self._gap_length: | ||
| return [] | ||
|
|
||
| gaps = [] | ||
|
|
||
| # First, collect all span start positions ('10' tokens) in the external range | ||
| span_starts = [] | ||
| for i, token_id in enumerate(request.prompt_token_ids): | ||
| if token_id == 10 and external_start <= i < external_end: | ||
| span_starts.append(i) | ||
| print(f"Found span starts at: {span_starts}") | ||
|
|
||
| # Create gaps for each span, bounded by the next span start | ||
| for idx, gap_start in enumerate(span_starts): | ||
| # Find the end of this span: either the next '10' or external_end | ||
| if idx + 1 < len(span_starts): | ||
| next_span_start = span_starts[idx + 1] | ||
| else: | ||
| next_span_start = external_end | ||
|
|
||
| span_length = next_span_start - gap_start | ||
| print( | ||
| f"Span at {gap_start}: length={span_length}, next_span at {next_span_start}" | ||
| ) | ||
|
|
||
| # Gap is min(self._gap_length, span_length), but not exceeding external_end | ||
| gap_end = min(gap_start + self._gap_length, next_span_start, external_end) | ||
|
|
||
| # Only add if we have at least some gap | ||
| if gap_end > gap_start: | ||
| print(f" Adding gap: ({gap_start}, {gap_end})") | ||
| gaps.append((gap_start, gap_end)) | ||
|
|
||
| print(f"Gaps: {gaps}") |
There was a problem hiding this comment.
Debug print statements should not be left in production code. All these print statements (lines 93-129, and elsewhere in this file) appear to be debugging output and should be removed or converted to proper logging.
|
|
||
| # Run the model. | ||
| # Use persistent buffers for CUDA graphs. | ||
| print(f"Putting positions: {positions}.") |
There was a problem hiding this comment.
Debug print statement should not be left in production code. This should be removed or converted to proper logging.
| print(f"Putting positions: {positions}.") | |
| logger.debug("Putting positions: %s.", positions) |
|
|
||
| if failed_kv_load_req_ids and not self.recompute_kv_load_failures: | ||
| requests = [self.requests[req_id] for req_id in failed_kv_load_req_ids] | ||
| print(f"Inserted {len(failed_kv_load_req_ids)} failed_kv_load_req_ids.") |
There was a problem hiding this comment.
Debug print statement should not be left in production code. This should be removed or converted to proper logging.
| print(f"Inserted {len(failed_kv_load_req_ids)} failed_kv_load_req_ids.") | |
| logger.warning( | |
| "Inserted %d failed_kv_load_req_ids.", | |
| len(failed_kv_load_req_ids), | |
| ) |
| dedup_bl = list({block.block_id: block for | ||
| block in blocks_list}.values()) |
There was a problem hiding this comment.
The variable name 'dedup_bl' is cryptic and should be renamed to something more descriptive like 'deduplicated_blocks' or 'unique_blocks' for better code readability.
| query_offset_1 = kv_head_idx * num_queries_per_kv + offs_m % num_queries_per_kv | ||
| query_offset = ( | ||
| offs_d_new = tl.arange(0, HEAD_SIZE_PADDED // 2) | ||
|
|
There was a problem hiding this comment.
There is inconsistent trailing whitespace after this line. Trailing whitespace should be removed for code cleanliness.
| def align_to_block_size(num_tokens: int, block_size) -> int: | ||
| """Align the number of tokens to the block size.""" | ||
| return (num_tokens - 1) // block_size * block_size | ||
| return (num_tokens // block_size) * block_size |
There was a problem hiding this comment.
There is inconsistent trailing whitespace after this line. Trailing whitespace should be removed for code cleanliness.
| return (num_tokens // block_size) * block_size | |
| return (num_tokens // block_size) * block_size |
|
|
||
| if not hasattr(self, "rotate"): | ||
| if not isinstance(self.model.model.layers[0], PPMissingLayer): | ||
| self.rotate = self.model.model.layers[0].self_attn.rotary_emb |
There was a problem hiding this comment.
Hey, this code assumes every model layer has a self_attn attribute, but that's not universal. GPT-OSS for example uses TransformerBlock which exposes attention as .attn instead of .self_attn, so this crashes with AttributeError: 'TransformerBlock' object has no attribute 'self_attn' on the first inference request.
I reproduced it with spnl-llm-d-cuda:v0.19.0 + GPT-OSS-120B. Suggested fix: replace the direct .self_attn access with a helper that checks both naming conventions:
def _get_attn_module(layer):
"""Get attention module from a layer, handling different naming conventions."""
for attr in ("self_attn", "attn"):
if hasattr(layer, attr):
return getattr(layer, attr)
raise AttributeError(
f"{type(layer).__name__} has no recognized attention attribute "
f"(tried 'self_attn', 'attn')"
)Then the existing code becomes:
if not hasattr(self, "rotate"):
if not isinstance(self.model.model.layers[0], PPMissingLayer):
self.rotate = _get_attn_module(self.model.model.layers[0]).rotary_emb
else:
for lay in self.model.model.layers:
if not isinstance(lay, PPMissingLayer):
self.rotate = _get_attn_module(lay).rotary_emb
break…stic. Signed-off-by: omerpaz95 <omerpaz95@gmail.com>
Signed-off-by: omerpaz95 <omerpaz95@gmail.com>
Signed-off-by: omerpaz95 <omerpaz95@gmail.com>
…3058) Signed-off-by: ramos <49182011+nemoramo@users.noreply.github.com> Signed-off-by: mayufeng <mayufeng@example.com> Co-authored-by: mayufeng <mayufeng@example.com>
de5f26d to
487794c
Compare
Make block size, gap policy, gap length, and pad token configurable via VLLM_V1_SPANS_* environment variables so they can be set in Docker deployments without CLI args. Also register the missing --gap-policy-name and --gap-policy-config CLI arguments that existed as EngineArgs fields but were not exposed to argparse. New env vars: - VLLM_V1_SPANS_PAD_TOKEN - VLLM_V1_SPANS_BLOCK_SIZE - VLLM_V1_SPANS_GAP_POLICY_ENABLE - VLLM_V1_SPANS_GAP_LENGTH
Add env var configuration for span/gap-policy parameters
| ), | ||
| # for block-attention, the token used for padding sequences | ||
| # to block boundaries (client-side) | ||
| "VLLM_V1_SPANS_PAD_TOKEN": lambda: int( |
…o-end counterparts for PIC chunk hashing and span boundary behavior
…solete test cases for improved clarity and maintainability (the chunk's 2 blocks are marked PIC)
…larity and update related test cases to improve prefix cache handling
…ity and assertions for baseline and marked requests
… prompt size and improve assertions for baseline and marked requests and create test_repeated_pic_span_reuse_and_gap_recompute_e2e
…ad the physical KV block bytes directly)
…ts chain-hashed block 2 ≠ the warmup's NONE_HASH-rooted block 2
…refix reuse and gap policy assertions
…in favor of generate_single_output
…ork according to per occurance occ1_kv and occ2_kv instead of shared block
…st_large_gap_length_does_not_livelock_e2e
…unwarmed_pic_chunk_halts_prefix_cache_reuse_e2e should have a comparison with independent warmup
Add Legolink tests + new test 4: PIC spans preserve prefix caching across requests
The render server already does GPU-less preprocessing (messages -> token_ids).
Add the symmetric postprocessing half: POST raw model output text and get back
structured {reasoning_content, content, tool_calls} using the server's
configured reasoning + tool-call parsers.
This lets a client that generates via raw /v1/completions (e.g. a spans
middleware that must control token_ids + span_starts) obtain the same
structured response /v1/chat/completions produces, without re-implementing the
MiniMax (or any) parsers.
- OpenAIServingRender: accept reasoning_parser; build reasoning_parser_cls;
add parse_chat_output (split reasoning, then extract tool calls from the
reasoning-stripped content, matching chat serving).
- Wire reasoning_parser through both init sites (generate router + CPU render
server).
- ParseRequest/ParseResponse models + POST /v1/chat/completions/parse route.
- Unit tests for the handler orchestration.
Signed-off-by: Itay Etelis <itay.etelis@ibm.com>
chore: remove spans redundant files
It is a GPU-less unit test of the render server's /parse handler, not an e2e or spans test; colocate it with the other render server tests.
…oint feat(render): add /v1/chat/completions/parse postprocessing endpoint
to be the same as that for /v1/chat/completions Signed-off-by: Doron Chen <doronchen@dhcp-9-147-172-114.givatayim.il.ibm.com>
Fix/parse marker gate
Creating a PoC/demo for using Legolink with Spans, using https://github.com/sdavidbd/vllm/tree/feature/segmented-prefill/ for creating virtual requests for the Legolink-required 32 tokens in the beginning of each span.