Conversation
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…orkspace-sizing-20260626 # Conflicts: # vllm/v1/worker/gpu_model_runner.py Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…0824 # Conflicts: # vllm/v1/attention/backends/flashinfer.py Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…0824 Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
The DCP prefill wrapper owns no workspace: it holds a paged and a ragged FlashInfer wrapper that captured the arena when it was built. It was never registered with the workspace state, and registering it would have rebound nothing anyway, so a later arena growth left both inner wrappers planning against the workspace they started with. Register the members instead, and run the registration for the DCP branch too. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com> (cherry picked from commit 07600ab)
capture_model() locked the workspace as its last step, but it returns early when both capture modes are disabled and the worker skips it entirely under enforce_eager. Those configurations started serving with the workspace still growable, so an attention wrapper could take it past the size KV sizing was told to expect. Lock it where warmup ends instead, which every configuration reaches. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com> (cherry picked from commit 9d2d73a)
…0824 Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…fore The lock sat ahead of the V1 sampler warmup, which runs a real dummy step and can be the first caller to need a larger attention workspace. Under enforce_eager nothing had locked it before, so that step would now hit a locked manager and abort startup. Move the lock past both late warmup branches, and assert the ordering in the regression test. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com> (cherry picked from commit ef19609)
Retire PR #108 from integration PR #115 with upstream PR vllm-project#47637. Restore upstream Triton specialization and remove the policy-only test. Keep the existing interpolation correctness tests unchanged. Co-authored-by: OpenAI Codex <codex@openai.com> Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…0824 Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…0824 Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…0824 Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
Resolve conflicts against upstream renames and deprecations: - use_dedicated_xqa -> use_xqa (XQA decode support on SM90) - drop _seq_lens_cpu / _num_computed_tokens_cpu from the ubatch CommonAttentionMetadata construction (removed in the 0.29 deprecation sweep); is_prefilling and mm_req_doc_ranges are now upstream fields and are still re-keyed per ubatch slice here - fold the DCP KV-sharing guard into the elif branch that follows the mm-prefix dispatch
…0824 # Conflicts: # vllm/v1/attention/backends/flashinfer.py # vllm/v1/worker/gpu/model_runner.py # vllm/v1/worker/ubatch_utils.py Resolutions: adopt upstream use_dedicated_xqa -> use_xqa rename and the KV-sharing DCP guard; drop the _seq_lens_cpu/_num_computed_tokens_cpu propagation upstream removed; keep the mm-prefix dispatch and the NVFP4-aware DCP cache view. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…0824 Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
The mask-owning mm-prefix prefill wrapper is built on the first batch that
carries bidirectional ranges, which is after execution has locked the
workspace arena, and it asks for the full default float workspace. The
locked arena cannot grow, so the first multimodal batch aborts the engine:
AssertionError: Workspace is locked but allocation from
'flashinfer.py:_allocate_workspace_buffer' requires 512.00 MB,
current size is 95.69 MB.
The profiling gate that was supposed to keep mm-prefix models off this path
does not prevent it: lock_workspace() is called unconditionally by the CUDA
graph capture and by the worker, so the gate only skipped the reservation
while the lock still happened.
Reserve the wrapper instead of bailing out:
- Drop the is_mm_prefix_lm gate from the persistent profiling support query.
- Grow the shared float arena to its default before the runtime wrappers are
built, so none of them can be the allocation that needs a locked arena to
grow. The arena only ever grows, so the final size is unchanged.
- Materialize and register the mm-prefix wrapper during profiling, gated on
the model and on the range buffers exactly as build() gates its use, so a
plain model never builds the JIT variant.
- Count the wrapper in _iter_workspace_wrappers() and in the reserve debug
info, so its int workspace is part of the measured reservation.
Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…0824 Both sides add tests to tests/v1/worker/test_attn_utils.py under the same @pytest.mark.parametrize opener: upstream's HiSparse block-size resolution and this branch's NVFP4 spec view-shape case. Keep both, each with its own decorator, and take the union of the imports they need. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
The rejection reads as though SM100 itself were the problem. It is not: the mask-owning variant needs the fa2 prefill kernels, stock fa2/fa3 cannot read an NVFP4 cache, and this backend therefore pins the wrapper to trtllm-gen whenever the cache is NVFP4 -- which in turn cannot run a custom variant. The two requirements exclude each other wherever an NVFP4 cache is selectable at all, which upstream means SM100. Comment and message only; the gate behaves exactly as before. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
The previous wording claimed the fa2 kernels cannot read an NVFP4 cache. That is wrong for the variant this PR builds: it detects the packed FP4 KV dtype, declares the per-block scale factors as additional tensors, and the forward path hands them over as kv_cache_sf; a conditional GPU test covers exactly that path. What is actually missing is a selectable end-to-end configuration. Upstream the NVFP4 cache dtype validates only on SM100, where trtllm-gen serves it, and trtllm-gen cannot run a custom variant. The combination of this variant with the upstream SM100 KV update and layout is unvalidated, so the gate stays conservative for that reason rather than a kernel-capability one. Docstring and message only; the gate behaves exactly as before. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…0824 Upstream now types key/value as optional for KV-sharing decoder layers. The NVFP4 fa2 prefill path indexes them directly, so it fails the same way DCP prefill does and now raises NotImplementedError instead of a type error. Same guard as on the NVFP4 KV cache branch. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…-20260824 Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
The mask this path needs is now packaged in FlashInfer, so the variant declaration and the JIT ABI no longer belong here. Drop the embedded VllmMMPrefixAttention CUDA struct, the URI and additional tensor/scalar argument builder, the packed-fp4 tensor-name special case, and the direct gen_customize_batch_prefill_module call, and construct FlashInfer's wrapper instead. The wrapper is resolved by name at call time rather than imported at module scope: a FlashInfer without it must demote this backend before selection, not raise ImportError for every FlashInfer user. The capability probe builds the module once through the public generator, so the behaviour of supports_mm_prefix and the per-layer-group re-probe is unchanged. Two call-site contracts change with the public API. The range tensor is handed over as (N, 2) rows instead of the flattened view the old ABI read, which the wrapper requires to be contiguous -- a row slice of the row-major buffer is. And q/k/v scales are passed as run() keywords instead of folding q_scale * k_scale into an sm_scale scalar by hand, which makes this path identical to the plain prefill wrapper alongside it. The kernel-level NVFP4 test moves to FlashInfer with the variant. The tests kept here cover what is still vLLM's: range construction and clipping, buffer reuse, backend eligibility and fail-closed startup, text-only and decode-only gating, short-extend reclassification, kernel block sizes, and an end-to-end paged-KV comparison against a dense reference. Also fix the range-splitting test to read seq_lens_cpu_upper_bound, the name the production call site already uses; CommonAttentionMetadata has no seq_lens_cpu. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
The composite multimodal-prefix backend landed upstream and appended its tests to the same file this branch appends to, so the two blocks collided on their shared boundary rather than on any shared code. Keep both: the composite tests and the FlashInfer adapter tests cover different backends and share no names. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
The composite reaches an mm-prefix batch by routing the bidirectional spans to Triton and the causal remainder to FlashInfer. FlashInfer serves the whole batch in one kernel when its own combination validates, so on SM100 it is offered first and the composite stays the fallback for what it rejects -- an fp8 cache, a trtllm-only page size, a build whose JIT variant will not compile. Selection already works this way: the priority list is walked in order and the first backend that validates wins, so nothing is forced. `FlashInferBackend.validate_configuration()` overrides the base method, which gained `use_rswa`; without it the override rejects the keyword and every selection through this backend raises TypeError. The two selector tests that enumerate the mm-prefix fallback order encoded the assumption that FlashInfer cannot serve these batches. It can once the wrapper is available, so both now take that as a parameter and assert the order it implies: FlashInfer when it validates, the composite or the Triton fallback when it does not. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
`blackwell_selection` and `hopper_selection` patched `vllm.v1.attention.backends.fa_utils.is_fa_version_supported`, but `get_flash_attn_version()` imports that symbol from `vllm.vllm_flash_attn.flash_attn_interface` inside the call, so the patch never applied and the fixtures read whatever the host actually supports. `_blackwell()` in the same file already patches the defining module; use that path in both fixtures. Also assert that the FlashInfer override forwards `use_rswa`. The override exists to reject mm-prefix with DCP; R-SWA is decided by the base class, so a dropped flag would silently accept a configuration the base would refuse, and a stale signature raises TypeError for every selection through this backend. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
FlashInfer can reserve its attention workspace after memory profiling has already sized the KV cache. That memory looks free while the cache is sized and is taken back on the first real request, so the engine can die on a batch it was told it had room for. Move the reservation into profiling. Wrappers on participating routes are materialized while measurement is running, their allocator use is visible to KV sizing, and the shared arena is locked after the sizing warmups so nothing below can grow it afterwards. Backends opt in through `PersistentWorkspaceProfilingSupport`; the fail-closed default stays `UNSUPPORTED`, so a backend that has not been audited keeps today's behaviour. Which FlashInfer routes need a persistent workspace is now resolved once per builder from the worst-case shape rather than per batch, because the decision has to be known before profiling runs, not after the first batch arrives. Composite backends are part of this. `TRITON_FLASHINFER` wraps a FlashInfer child that does reserve a workspace, but the composite builder forwarded only the buffer getter and setter, so it inherited the fail-closed default and the lifecycle was skipped for it -- while the arena was still locked afterwards. The composite now composes its children's support and forwards all three reservation calls to both of them: either child can be routed to at runtime, one `UNSUPPORTED` child disables the pair, and one `REQUIRED` child makes the pair `REQUIRED`. `TritonAttentionMetadataBuilder` declares `NEUTRAL`, which it has to do explicitly -- it holds no workspace of its own, but the fail-closed default would otherwise disable the lifecycle for every composite wrapping it. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…site The runner prefers `get_workspace_buffer_state()` over the bare buffer because the state carries the registered wrappers as well as the arena: one state tracks every wrapper so they can all be rebound together when the arena grows. The composite forwarded only the buffer, so each child kept its own state and a wrapper registered on one would not be rebound when the other grew the arena. Forward the state API, and keep the bare buffer as the fallback for a backend that does not implement it. Compose the children's profiling support fail-closed as well. Anything that is not one of the three declared values now yields `UNSUPPORTED`: previously an invalid return was laundered into `NEUTRAL` or `REQUIRED`, hiding it from the gate that would otherwise have caught it. Also correct why `TritonAttentionMetadataBuilder` is `NEUTRAL`. It does hold persistent tensors -- the softmax segment buffers and the R-SWA prefix lengths -- but they are constructor-owned and kept alive by the lease, with nothing reserved lazily and nothing to rebind. Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
# Conflicts: # tests/v1/attention/test_mm_prefix.py Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…-cherry-20260914 # Conflicts: # tests/v1/attention/test_mm_prefix.py # vllm/v1/attention/backends/flashinfer.py Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…-cherry-20260914 # Conflicts: # tests/v1/worker/test_cudagraph_memory_profiling.py # vllm/v1/attention/backends/flashinfer.py Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
…-20260824 Signed-off-by: lesj0610 <lesj0610@users.noreply.github.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Purpose
Fresh integration branch replacing the previous integration stack (#99), rebuilt from current upstream
mainwith only the up-to-date revisions of the surviving fork PRs. Closed/superseded work (#101, #107) and the withdrawn Qwen3-VL specialization override (#108 / vllm-project#47637) are excluded, and the three FlashInfer-stack PRs are integrated in their rebuilt forms.Composition
Merged on top of upstream main (
0ecc284790), in order:lesj/cutlass-fp8-capability-gate-20260822lesj/fused-gdn-decode-ops-guard-20260816lesj/unify-page-scale-block-20260822agent/mrope-cache-position-boundsagent/qwen35-processor-backenduse_fastfixlesj/gemma4-flashinfer-selectorlesj/flashinfer-nvfp4-kv-cache-20260624lesj/use-flashinfer-workspace-sizing-20260626lesj/flashinfer-mm-prefixIntegration notes
cudagraph_utilswhile [Core] Profile persistent FlashInfer workspace before KV sizing #102 grew the same entry points on the runner. The follow-up commit grafts [Core] Profile persistent FlashInfer workspace before KV sizing #102's persistent-workspace accounting onto main's flow: the runner keeps [Core] Profile persistent FlashInfer workspace before KV sizing #102's method contract,profile_cudagraph_memoryreserves the attention workspace inside the cleanup-guarded span, reports it oncudagraph_memory_persistent_estimate, and returns main's graph-memory estimate.o_data_typeselection; restored.use_dedicated_xqa/model_stateattributes) and were fixed here.tests/kernels/attention/test_flashinfer.pyreferenceFlashInferBackend.get_kv_cache_shape, which no longer exists on the [Core] Use FlashInfer for pre-SM100 NVFP4 KV cache updates #100 branch or upstream main; the 35 functional nvfp4 tests pass.Test Result
On one SM80 device, FlashInfer v0.6.17 +
variant_owns_mask(editable):ruff check/ruff formatand commit hooks (incl.mypy) passed on all touched filesRemoval of #108 / vllm-project#47637 (2026-09-05)
Restore the bilinear position-embedding kernel's default
@triton.jitspecialization and remove the dedicated specialization-policy assertion. The existing interpolation correctness tests remain unchanged.Both affected files exactly match the upstream revision already integrated into this branch. Shape-dependent JIT compilation follows the upstream policy again.
CUDA validation on GPU 1 (NVIDIA CMP 170HX, SM80), Python 3.12.12, PyTorch 2.13.0+cu130, Triton 3.7.1:
Result: 18 passed. Syntax compilation,
git diff --check, and all applicable pre-commit hooks passed for both affected files.Serving/model end-to-end evaluation and first-request JIT latency were not measured for this removal.
AI assistance: Codex, Claude