Skip to content

[POC]Segmented spans#2

Open
omerpaz95 wants to merge 68 commits into
mainfrom
segmented_spans
Open

[POC]Segmented spans#2
omerpaz95 wants to merge 68 commits into
mainfrom
segmented_spans

Conversation

@omerpaz95

Copy link
Copy Markdown
Owner

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.

David Ben-David and others added 4 commits January 21, 2026 17:00
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>
@omerpaz95 omerpaz95 self-assigned this Feb 10, 2026
@github-actions

Copy link
Copy Markdown

👋 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 fastcheck CI which starts running only a small and essential subset of CI tests to quickly catch errors.

You ask your reviewers to trigger select CI tests on top of fastcheck CI.

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 ready label to the PR or enable auto-merge.

If you have any questions, please reach out to us on Slack at https://slack.vllm.ai.

🚀

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Comment thread vllm/v1/core/kv_cache_utils.py Outdated
Comment on lines +556 to +557
firstok = curr_block_token_ids[0]
if firstok == envs.VLLM_V1_SPANS_TOKEN_PLUS:

Copilot AI Feb 23, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The variable name 'firstok' appears to be a typo for 'first_tok' or 'first_token'. This should be renamed for clarity.

Suggested change
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:

Copilot uses AI. Check for mistakes.
Comment thread vllm/v1/core/kv_cache_utils.py Outdated
Comment on lines +623 to +626

extra_keys = recompute_token_handler(
block_tokens[0], block_tokens[:start_token_idx], extra_keys)

Copilot AI Feb 23, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There is inconsistent trailing whitespace after this line. Trailing whitespace should be removed for code cleanliness.

Suggested change
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
)

Copilot uses AI. Check for mistakes.
mask=dim_mask_b[None, :] & query_mask_0[:, None] & query_mask_1[:, None],
other=0.0,
)

Copilot AI Feb 23, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There is inconsistent trailing whitespace after this line. Trailing whitespace should be removed for code cleanliness.

Suggested change

Copilot uses AI. Check for mistakes.
Comment thread vllm/envs.py Outdated
VLLM_LOG_MODEL_INSPECTION: bool = False
VLLM_DEBUG_MFU_METRICS: bool = False


Copilot AI Feb 23, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There is inconsistent trailing whitespace after this line. Trailing whitespace should be removed for code cleanliness.

Suggested change

Copilot uses AI. Check for mistakes.
Comment on lines +93 to +129
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}")

Copilot AI Feb 23, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copilot uses AI. Check for mistakes.
Comment thread vllm/v1/worker/gpu_model_runner.py Outdated

# Run the model.
# Use persistent buffers for CUDA graphs.
print(f"Putting positions: {positions}.")

Copilot AI Feb 23, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Debug print statement should not be left in production code. This should be removed or converted to proper logging.

Suggested change
print(f"Putting positions: {positions}.")
logger.debug("Putting positions: %s.", positions)

Copilot uses AI. Check for mistakes.

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.")

Copilot AI Feb 23, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Debug print statement should not be left in production code. This should be removed or converted to proper logging.

Suggested change
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),
)

Copilot uses AI. Check for mistakes.
Comment thread vllm/v1/core/block_pool.py Outdated
Comment on lines +401 to +402
dedup_bl = list({block.block_id: block for
block in blocks_list}.values())

Copilot AI Feb 23, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copilot uses AI. Check for mistakes.
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)

Copilot AI Feb 23, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There is inconsistent trailing whitespace after this line. Trailing whitespace should be removed for code cleanliness.

Suggested change

Copilot uses AI. Check for mistakes.
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

Copilot AI Feb 23, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There is inconsistent trailing whitespace after this line. Trailing whitespace should be removed for code cleanliness.

Suggested change
return (num_tokens // block_size) * block_size
return (num_tokens // block_size) * block_size

Copilot uses AI. Check for mistakes.
Comment thread vllm/v1/worker/gpu_model_runner.py Outdated

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

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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>
omerpaz95 pushed a commit that referenced this pull request Feb 25, 2026
…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>
omerpaz95 and others added 4 commits March 16, 2026 11:34
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
Comment thread vllm/envs.py
),
# for block-attention, the token used for padding sequences
# to block boundaries (client-side)
"VLLM_V1_SPANS_PAD_TOKEN": lambda: int(

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

this isn't used

almogtavor and others added 30 commits May 12, 2026 21:02
…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
…ts chain-hashed block 2 ≠ the warmup's NONE_HASH-rooted block 2
…ork according to per occurance occ1_kv and occ2_kv instead of shared block
…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>
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>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants