Fix a bug in 1D input shape by WoosukKwon · Pull Request #5 · vllm-project/vllm

WoosukKwon · 2023-03-03T04:18:38Z

This PR fixes a miscalculation of the input shape when iteration-level scheduling is used.

More improvements awq

* finish changing scheduler * finish merge * fix model * Fix (vllm-project#5) * fix problems * fix * delete unused params * remove redundant comments --------- Co-authored-by: Xiangyu Tian <109123695+xiangyuT@users.noreply.github.com>

Align optimum-intel based model signature with vLLM signature

…imum Install optimum-intel from latest main

* Drop indecies when finish * min 1 attention layer * CG is working on forward pass passing * Remove comments * cosmetics - rename indecies -> indices, organize some whitespaces * Add some TODOs * Adding mamba cache for cg * Remove useless vars from input_metadata * Remove unused import * Set the seqlen offset to boolean * Return only hidden state * Return only hidden states * Add padding to match forward pass bs * Is prompt instead of seqlen offset * Remove mamba cache class (not used) * Another remove * Remove * Use mamba4gc * Fix mamba forward, run update only on non prompt * Use 1 index after the maximal index * Remove import * Remove import * typo * typo * place holder * Padding and empty token takes it from the first empty place * reformat * Apply suggestions from code review Whitespaces --------- Co-authored-by: Mor Zusman <morz@ai21.com> Co-authored-by: Tomer Asida <tomera@ai21.com> Co-authored-by: tomeras91 <57313761+tomeras91@users.noreply.github.com>

…3small [Model][Kernels] Support Phi3small architecture, blocksparse attnention prefilling kernel, CUDA+Triton paged attn kernels

Faster v2 hopper fused moe kernel configs

Kuntai disagg refactor

…3 + walker nuke + bug-class-vllm-project#1 ratchet Establishes the typed `MegaTape<S>` lowering pipeline in a new `ferrite-mega-ir` crate and migrates the megakernel codegen off the walker-based `emit_op` design onto a typed IR that the proc-macro constructs at expansion time. Per-op substrate witnesses move runtime checks (page-slot bounds, lifecycle, mbarrier phase math, scratch overlap, scratch budget, warp roles) toward construction- time invariants — the plan's six bug classes. Sprint 0+1a+1b — substrate vocabulary + crate extraction: - `ferrite-mega-ir` crate carved out of the macro crate so the IR is a real type-level artifact, not a TokenStream stub. - Sealed substrate witnesses: `Page<ID, S, State>` with lifecycle states (`Empty`/`Filled`/`Produced`), `ScratchRegion<OFFSET, BYTES, Scope>`, `MbarrierPhase<P>`, `WarpRoleTag<R>`, `IsValidSlot`, `Disjoint`, `WithinBudget`, `ArriveCountToPhase`. - MegaIR decoupled from `W` (the per-canonical Weights type) so the IR can be constructed without a model in scope. Sprint 2 — gate-up + qkv-rope + downproj migrations: - silu_upgate, gelu_upgate, down_proj_residual, fused_qkv_rope_cache lower through the typed IR. Each variant's `emit_*` function takes a typed node and renders CUDA via `ferrite_mega_ir::emit`, replacing the walker's per-variant bespoke `emit_op` arm. Sprint 3 — fused_add_rms_norm + embed migrations. Walker nuke: - 19 per-variant `emit_X` walker fns deleted; the typed pipeline is now the only path. `EmitCtx` + `WalkerLines` + the `emit_op_from_typed` bridge are gone — `ferrite_mega_ir::emit` is the single render surface. - `BaseStage::new(n, op_page_count, num_pages)` validates page budget at IR construction (bug class vllm-project#1). Tautological in the proc-macro today (base_stage always 0), but the typed gate is load-bearing for Phase 4 cross-op pipelining where non-zero `base_stage` will fail at construction instead of corrupting emitted CUDA at runtime. Squashes 123 in-progress commits (Sprints 0+1+2+3 iters, walker deletion, BaseStage ratchet) into one landing commit. Backup of the original history at `backup/worktree-ff-mega-codegen-pre-squash`. Plan: `MEGA_IR_PLAN.md` — Sprint A done (RmsNorm), Sprint B done (FusedQkvRopeCache), Sprint C in progress (FusedAddRmsNorm + SiluUpgate + GeluUpgate + DownProjResidual landed; remaining Sprint D ops pending). Bug class vllm-project#1 caught by typed gate; bug classes vllm-project#2/vllm-project#3/vllm-project#4 are NIGHTLY TODOs (need per- variant const generics on MegaOp values, blocked on stable-Rust limits noted in `BaseStage` type-level docs); bug class vllm-project#5 (scratch budget overrun summed across the tape) is the next ratchet. Signed-off-by: Nick Mitchell <nickm@us.ibm.com>

Per MEGA_IR_PLAN.md §10: substrate-aware lowering for the second Instruction variant. Bug classes targeted: vllm-project#1 (page bounds), vllm-project#2 (lifecycle: in / qkv-weight / cos-sin / q-out / k-out / v-out across loader → consumer → storer), vllm-project#3 (per-iter mbarrier phase math), vllm-project#4 (scratch overlap inside the per-tok loop), vllm-project#5 (within-budget scratch), vllm-project#6 (warp-role pairing). substrate.rs: - New `RopeScope` (sealed `IsScratchScope` + `IsScratchScopePub`) for the per-token Q/K rotation tiles. Sibling regions in this scope are required to be `disjoint_with`-discharged. - New `IterCount(u32)` newtype with `iters > 0` invariant — used to bump the cumulative arrive count by `iters * arrives_per_iter` (per-iter math is variant-internal: with arrives_per_iter == 1, the kernel's `phase ^= 1` toggle is correct iff the boundary phase matches the cumulative count parity at op start, which is what the lowering validates via MbarrierPhase::assert_matches). nodes.rs: - New `RotaryRef` helper newtype (non-empty path). Lets the typed node distinguish global rotary from per-layer-local rotary (`wm.rotary` vs `wm.rotary_local`) per the typed-fanout walker's injection. - New `MegaNode::FusedQkvRopeCache` variant. Substrate-proof fields: 6 PageIds (lifecycle-witnessed), 2 disjoint ScratchRegion<RopeScope> (Q/K rotation tiles), consumer_phase + storer_phase + iters, four WarpRoleTag<R>. Helper fields alongside: layer / qkv_weight / rotary / biased / interleaved. lower.rs: - `MegaTapeBuilder::push_fused_qkv_rope_cache` walks every page through Empty → Filled → Produced → Empty (typestate-burned), packs Q/K rotation tiles back-to-back in scratch, validates phase parity at the op boundary, releases pages, and bumps the cumulative arrive count by `iters` (per-iter math falls out of `(C+t)&1 == (C&1)^(t&1)`). - `lower()` arm dispatches `Instruction::FusedQkvRopeCache` with a per-tape `SlotAllocator` that picks 5 non-aliasing output / weight page ids modulo `num_pages`. Real pipeline-aware allocation lands in Sprint D. - `OpInput::iters: u32` added (default 1 via `OpInput::new`). weight_paths arity for Rope is required to be 2 (qkv-packed, rotary). New `LowerError::WrongWeightArity` and `LowerError::SubstrateBudgetTooSmall` variants. tape.rs: - `impl Debug for MegaTape` (size + substrate budget summary), so test panics can format `Result<MegaTape, LowerError>`. 9 new tests in `lower::tests` covering: well-formed Rope round-trip; arrive-count phase advance across an op chain (5-iter odd flips parity, 4-iter even keeps it); Q/K scratch overlap rejection; within-op page aliasing rejection; iters=0 rejection; empty-rotary-path rejection; lower() round-trip; weight-paths arity error. Existing RmsNorm tests updated for irrefutable-pattern → `let-else` since `MegaNode` now has 2 variants. 30 tests pass on H100 pod (was 21). `cargo clippy -p ferrite-mega-ir --all-features -- -D warnings` clean. No emit yet — Sprint A — D land the typed lowering for every variant first, syntactic emit after. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> Signed-off-by: Nick Mitchell <nickm@us.ibm.com>

…ilu,Gelu}Mul Per MEGA_IR_PLAN.md §10: substrate-aware lowering for the residual fold (`Instruction::Add` — what the plan calls "DownProjResidual"), the fused add+norm (pre-attention / pre-MLP layer entry), and the gate-up MLP fusion in both Silu and Gelu activation forms. Bug classes vllm-project#1–vllm-project#6 all exercised between the four variants. substrate.rs: - New `MlpScope` (sealed `IsScratchScope` + `IsScratchScopePub`) for the gate-up fusion's per-token-iter gate / up activation tiles. Sibling regions in this scope MUST be `disjoint_with`- discharged at lowering time (Sprint C's bug class vllm-project#4 site). nodes.rs: - New `MegaNode::Add`: 2 PageIds (delta + residual, lifecycle-walked), consumer/storer phase, 4 warp-role tags. No weight, no scratch — element-wise residual fold lives entirely in registers. - New `MegaNode::FusedAddRmsNorm`: 3 PageIds (delta + residual + weight) lifecycle-walked, partial_sums in `RmsNormScope` (sibling regions on a future RmsNorm op `disjoint_with`-discharged), phase parity, 4 role tags, layer + weight helpers. - New `MegaNode::FusedGateUpActivateMul` (single variant covering both `FusedGateUpSiluMul` and `FusedGateUpGeluMul` instructions — same substrate shape, distinguished by `activation: GateUpActivation { Silu | Gelu }`): 3 PageIds (in + packed gate-up weight + out), 2 disjoint `MlpScope` tiles (gate_buf, up_buf), boundary phase parity, IterCount, 4 role tags, layer + weight + activation helpers. lower.rs: - `MegaTapeBuilder::push_add`, `push_fused_add_rms_norm`, `push_fused_gate_up_activate_mul` — substrate-proof discharge per call: page allocation + bounds (vllm-project#1), lifecycle walks Empty → Filled → Produced → Empty (vllm-project#2), scratch (vllm-project#4 disjointness + vllm-project#5 within-budget), boundary phase parity (vllm-project#3 — per-iter math for the gate-up multi-iter shape falls out of `(C+t)&1 == (C&1)^(t&1)`), warp-role pinning (vllm-project#6). - `lower()` arms for `Instruction::Add`, `FusedAddRmsNorm`, `FusedGateUpSiluMul`, `FusedGateUpGeluMul`. - New `pick_distinct_slot` helper: picks the next free page id that doesn't collide with caller-supplied input/output slots. Used by Sprint C arms when the Instruction only supplies in/out and the lowering needs to synthesize a non-aliasing weight slot. 13 new tests in `lower::tests` covering: well-formed Add round-trip; self-aliasing rejection; arrive-count phase advance across two Adds; FusedAddRmsNorm well-formed shape; delta/residual aliasing rejection; gate-up Silu and Gelu lowering; gate/up scratch overlap rejection; Instruction-side round-trip for each variant; multi-op chain threading phase advance through Add → FusedAddRmsNorm → FusedGateUp (verifies the cumulative arrive count carries the correct parity at each boundary, including across the multi-iter gate-up bump). `lower_unmigrated_variant_returns_not_yet_lifted` test repointed at `Gemm` (Sprint D scope) since `Add` now lowers. 43 tests pass on H100 pod (was 30); `cargo clippy -p ferrite-mega-ir --all-features -- -D warnings` clean. No emit yet. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> Signed-off-by: Nick Mitchell <nickm@us.ibm.com>

… canonicals lift cleanly Wire the next-biggest unblocker. After this every bf16 llama-3.2-1B canonical emits a compile-time-substrate-proof-bearing `build_mega_tape_*()` fn. The user-build monomorphization fires every `const { assert!(...) }` block in mega_ir::nodes for the const args the proc-macro computed. ferrite-mega-ir/nodes.rs: - New `FusedCublasGemmAdd` MegaNode variant. Substrate shape: 3 pages (in / weight / residual=output where the gemm result is folded back in place), `GemmScope` B-tile, multi-iter, layer + matmul shape. Constructor's `const {}` block proves: page bounds (vllm-project#1), pairwise non-alias (vllm-project#2 within-op), scratch within-budget (vllm-project#5), iters > 0, n/k > 0, consumer/storer phase parity (vllm-project#3). ferrite-mega-ir/lower.rs: - New `MegaTapeBuilder::push_fused_cublas_gemm_add::<...>` builder method. 13 const generics (3 page ids + 2 scratch + 2 phases + iters + layer + n + k + num_layers + arrives). ferrite-forward-macro/codegen.rs: - Dispatch arm for `Instruction::FusedCublasGemmAdd(in_slot, residual_slot, layer, n, k)` synthesizes a non-aliasing weight slot, packs the B-tile across all of `SCRATCH_BYTES`, and emits the literal-const-arg builder call. Per-canonical status (llama-3.2-1B): - bf16 m=1 sk=128 / m=1 sk=192 / m=8 sk=128 / m=8 sk=256 / m=8 sk=512 / m=512 sk=256: **emits build_mega_tape_*** ✓ (6/6) - awq-gemm / ct-int4-sym: dispatch needed for MarlinFusedQkvRope* - bnb-nf4-dq: dispatch needed for Bnb4FusedQkvRope* - fp8-* / ggml: dispatch needed for Fp8/GgmlFusedQkvRope* `phase-c step 1 — 6 canonical(s) emitted, 0 skipped` for the bf16 variant of llama-3.2-1B. The 8 quantization-flavored variants need their respective Marlin / Bnb4 / Fp8 / Ggml QKV-rope cache MegaNode variants (Sprint E continuation). End-to-end on H100 pod: - cargo test -p ferrite-mega-ir: 32 lib + 1 compile-fail driver green. - FERRITE_MEGA=1 FERRITE_MODELS=llama-3.2-1b cargo build -p vllm-cli --features cuda --release: clean. Every emitted build_mega_tape_*() fn user-builds, every const-generic push monomorphizes, every const {} substrate-proof block discharges. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> Signed-off-by: Nick Mitchell <nickm@us.ibm.com>

…dation mode) + bisect knobs Implements the free-running draft's core move, env-gated default-off (VLLM_SELF_SPEC_AHEAD_CHAIN): after each propose, run_ahead_chain() CONTINUES the real chain K+1 steps on the side stream during the verify (input = its own last token; first ahead token = the bonus guess), with real position/metadata updates and real KV writes. VALIDATION mode: outputs feed a per-request hit-rate counter (hit = all-K accepted AND bonus == guess, the free-running prior ~beta^(K+1)) and are then discarded -- the normal propose still runs, so the served output is byte-identical while the overlap machinery runs for real. Supporting changes: - scheduler: num_lookahead_tokens K -> 2K+2 under the flag, so ahead KV writes (up to committed+2K+2) always land in the request's own allocated blocks. - metadata privatization: the ahead loop runs on a shallow-copied CommonAttentionMetadata with CLONED seq_lens + block_table_tensor. The lockstep chain safely mutates runner-owned views because everything is serial; free-running, those in-place advances race the runner's next-step prep + the verify (torn values -> FA3 reads garbage pages). First free-running-specific isolation rule (OV1 rule vllm-project#5). - runner boot-retry: EADDRINUSE races on this shared box (other users' engines grab ports) could wedge a worker 600 s on a dead port; the runner now detects failed boots early, kills, and retries up to 3x. Debug state (bisection so far, crash = sticky CUDA illegal memory access): - serialized ahead chain, KV on: CLEAN, accept 2.91 (loop intrinsically OK, incl. extended-lookahead KV writes) - serialized, KV suppressed: CLEAN - concurrent, full: CRASHES - concurrent, KV suppressed: CRASHES -> live KV writes exonerated Remaining suspects: live metadata update/build vs verify, LM-head sampling vs verify, live-input piece replay vs verify. Bisect knobs: W7_AHEAD_STEPS, W7_AHEAD_NO_KV, W7_AHEAD_NO_SAMPLE, W7_SHADOW_MAIN_STREAM. CUDA_LAUNCH_BLOCKING is unusable here (breaks multi-rank NCCL init). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

…-agnostic wave wake - Remove in-core DP-rank routing from async_llm.py (drop the ROCm auto-route block, _pick_dp_rank_for_request / _ensure_disagg_transfer_id helpers, and the hashlib/current_platform imports). Rank pinning is delegated to the llm-d sidecar/router, per njhill/NickLucche. Delete the now-obsolete tests/v1/engine/test_dp_rank_routing.py. - Reframe the DP first-wave wake in core.py from a current_platform.is_rocm() gate to the external-LB DP mode (data_parallel_external_lb). In external LB the router addresses engines directly, so the coordinator is not in the per-request wake path; internal/hybrid LB stays bit-identical to upstream. Add a debug trace at the wake point (njhill vllm-project#5). - Revert VLLM_ENGINE_HANDSHAKE_TIMEOUT_MINUTES; restore the constant HANDSHAKE_TIMEOUT_MINS = 5 (njhill vllm-project#6). - Guard against division by zero on data_parallel_size_local in moriio_common.py port allocation (Doug vllm-project#10). - Fill the AITER root-cause/fix link (ROCm/aiter#3658) in the sampler warm-up skip TODO. - Pin the decode leg to the prefill DP rank via the X-data-parallel-rank header in the moriio toy proxy example. Signed-off-by: Shiksha Patel <shikpate@amd.com>

Ten changes, each forced by a measurement rather than a preference: - sanity vllm-project#1's threshold moves from absolute error to cosine. K carries attention-sink outliers (|k| up to 314 against a mean of 2.35), so 1e-2 absolute sits below vLLM's own numerical noise and is unreachable. - tau is calibrated per layer, not globally: the recomputation noise floor spans 0.99999 at layer 0 down to 0.964 at layer 26. - tau must be calibrated under a *partial* prefix hit. A full re-prefill reproduces bit for bit, which would falsely report zero noise. - The sensitivity sweep drops 0.90/0.95: different-content K pairs already score 0.90-0.96, so those thresholds sit inside the false-positive baseline. - The ample tier is defined per workflow. One workflow needs 80k-135k tokens and ten of them exceed every card we have. - Limited capacity uses mean(W), floored so the pool can still hold the longest single call; a smaller pool deadlocks the scheduler silently. - Storage re-estimated from 90 GB to 444 GB; the original figure mistook a configuration's working set for a workflow's. - max_tokens=1024 and gpu_memory_utilization=0.95 keep both capacity tiers on one card type, which removes a GPU architecture from inside the Q3 join. - 4/8 tiers skip the dump and reuse the 16 tier's tensors. - sanity vllm-project#5 is rewritten: the prefix-hit-ratio probe stopped discriminating once the proposers took distinct system prompts. Adds the Athena execution plan: the cluster ships only cu130 wheels for this base commit while the nodes run 535/550 drivers, so vLLM has to be built from source against a locally installed CUDA 12.9. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

…) (vllm-project#5996) ### What this PR does / why we need it? **Scope of Changes**: | File Path | | :--- | | `.../distributed/kv_transfer/kv_pool/ascend_store/ascend_store_connector.py` | | `vllm_ascend/distributed/kv_transfer/kv_pool/ascend_store/backend/backend.py` | | ` .../distributed/kv_transfer/kv_pool/ascend_store/backend/memcache_backend.py` | | ` .../distributed/kv_transfer/kv_pool/ascend_store/backend/mooncake_backend.py` | | ` vllm_ascend/distributed/kv_transfer/kv_pool/ascend_store/config_data.py` | | ` vllm_ascend/distributed/kv_transfer/kv_pool/ascend_store/kv_transfer.py` | | ` vllm_ascend/distributed/kv_transfer/kv_pool/ascend_store/pool_scheduler.py` | | ` vllm_ascend/distributed/kv_transfer/kv_pool/ascend_store/pool_worker.py` | | ` .../distributed/kv_transfer/kv_pool/cpu_offload/cpu_kv_cache_manager.py` | | ` .../distributed/kv_transfer/kv_pool/cpu_offload/cpu_offload_connector.py` | | ` vllm_ascend/distributed/kv_transfer/kv_pool/cpu_offload/metadata.py` | | ` vllm_ascend/distributed/kv_transfer/kv_pool/ucm_connector.py` | | ` vllm_ascend/distributed/kv_transfer/utils/mooncake_transfer_engine.py` | | ` vllm_ascend/distributed/kv_transfer/utils/utils.py` | | ` vllm_ascend/kv_offload/cpu_npu.py` | | ` vllm_ascend/kv_offload/npu.py` | | ` vllm_ascend/lora/lora_ops.py` | | ` vllm_ascend/lora/punica_npu.py` | | ` vllm_ascend/lora/utils.py` | ### Does this PR introduce _any_ user-facing change? ### How was this patch tested? - vLLM version: v0.13.0 - vLLM main: vllm-project@2c24bc6 --------- Signed-off-by: MrZ20 <2609716663@qq.com> Signed-off-by: SILONG ZENG <2609716663@qq.com>

…lm-project#6747) ### What this PR does / why we need it? | File Path | | :--- | | `tests/e2e/singlecard/compile/backend.py` | | `tests/e2e/singlecard/compile/test_graphex_norm_quant_fusion.py` | | `tests/e2e/singlecard/compile/test_graphex_qknorm_rope_fusion.py` | | `tests/e2e/singlecard/compile/test_norm_quant_fusion.py` | | `tests/e2e/singlecard/model_runner_v2/test_basic.py` | | `tests/e2e/singlecard/test_aclgraph_accuracy.py` | | `tests/e2e/singlecard/test_aclgraph_batch_invariant.py` | | `tests/e2e/singlecard/test_aclgraph_mem.py` | | `tests/e2e/singlecard/test_async_scheduling.py` | | `tests/e2e/singlecard/test_auto_fit_max_mode_len.py` | | `tests/e2e/singlecard/test_batch_invariant.py` | | `tests/e2e/singlecard/test_camem.py` | | `tests/e2e/singlecard/test_completion_with_prompt_embeds.py` | | `tests/e2e/singlecard/test_cpu_offloading.py` | | `tests/e2e/singlecard/test_guided_decoding.py` | | `tests/e2e/singlecard/test_ilama_lora.py` | | `tests/e2e/singlecard/test_llama32_lora.py` | | `tests/e2e/singlecard/test_models.py` | | `tests/e2e/singlecard/test_multistream_overlap_shared_expert.py` | | `tests/e2e/singlecard/test_quantization.py` | | `tests/e2e/singlecard/test_qwen3_multi_loras.py` | | `tests/e2e/singlecard/test_sampler.py` | | `tests/e2e/singlecard/test_vlm.py` | | `tests/e2e/singlecard/test_xlite.py` | | `tests/e2e/singlecard/utils.py` | ### Does this PR introduce _any_ user-facing change? ### How was this patch tested? - vLLM version: v0.15.0 - vLLM main: vllm-project@9562912 --------- Signed-off-by: MrZ20 <2609716663@qq.com>

Implemented and measured on RTX 5090 with real EXL3 Viterbi trellis: - vllm-project#6 Hessian-weighted scale: no gain (synthetic Hessian too uniform) - vllm-project#2 Adaptive lattice: no gain (Lloyd-Max already optimal for Gaussian) - vllm-project#3 Low-rank subspace: no gain (Hadamard makes residuals full-rank) - vllm-project#5 Sparse residual: significant — K3+sparse_only0.5% beats K3 at 3.2 bits - vllm-project#7 Multi-codebook (AQLM): fails (codebooks can't capture i.i.d. Gaussian) - #1 Matryoshka approx: no gain (single step insufficient) Pareto frontier: K4 standalone still optimal at 4 bits. Residual approaches win at non-standard bit widths (3.2, 5, 5.2, 7 bits). Sparse+Lloyd-Max combo gives 95.3% gap at 4.04 bits vs 94.8% at 4.0 bits. Co-authored-by: Claude <noreply@anthropic.com>

The fused tilelang prenorm kernels re-read `fn` from every token tile, which makes them L2-bound at prefill shapes (the best swept config (8, 6) runs 442 us at T=2048, ~155 GB/s). The operation is just a [T, 16384] bf16 x [24, 16384] GEMM plus a per-row fp32 sqrsum, and unlike the fused kernel the two-pass split reads x exactly twice: cuBLAS takes the GEMM (bf16 out, fp32 accumulate, upcast into the fp32 buffer -- torch.mm will not emit fp32 from bf16 in) and a new one-pass Triton reduction (_row_sqrsum_kernel) takes the sqrsum. The bf16 `fn` copy is cached on the weight tensor itself. Measured (harness `--kernel prenorm-gemm`, new cublas arm), vs tilelang (8, 6): 17.4 vs 26.2 us at T=64, 17.5 vs 26.9 at 128, 23.0 vs 82.6 at 512, 35.1 vs 201.6 at 1024, 85.0 vs 442.2 at T=2048 -- 5.2x at the prefill chunk shape, beating the >= 3x acceptance and the 95-130 us projection. Timing the production router end-to-end reproduces the direct-kernel numbers (6.7 / 16.7 / 85.1 us at T = 1 / 32 / 2048). The named risk (cuBLAS picking a degenerate tile at N=24) did not materialize: 808 GB/s on minimal bytes = ~1.6 TB/s on actual traffic, at the lm_head anchor rate. The route runs under do_bench_cudagraph in the harness, so cudagraph capture is exercised. Routing: the small-T one-CTA-per-token kernel keeps T < 32 (7.0 vs 10.5 us at T=1; T=32 is a dead heat at 16.8 vs 16.7), so _PRENORM_SMALL_T moves 64 -> 32 -- the re-measured crossover against the new backend. The block_m (8, 6) route stays in-tree behind _PRENORM_USE_CUBLAS as the escape, still tested and benched. Decode at batch 1 is untouched. Numerics: `fn` is rounded to bf16 (~3 mantissa bits below the fp32 reference) and the GEMM result is rounded to bf16 before the upcast. The parity test now pins this directly on `out` at rel 5e-3 against the tensor scale for every routed T -- the downstream mhc_pre tolerance alone cannot distinguish bf16-fn rounding from a broken kernel -- while T < 32 keeps the strict fp32 tolerance (atol 1e-5, rtol 1e-4). New boundary cases at T in {31, 32}; suite 59 -> 65, all passing, including the fuse_norm downstream cases which now exercise the cuBLAS route at T=128. Expected end-to-end: the prenorm GEMM is 16.5% of 32k prefill / 12.7% at 128k, so ~9-13% TTFT at 32k from a 5.2x kernel win. Server confirmation belongs to batch run vllm-project#5. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: haosdent <haosdent@gmail.com>

serve.sh -v 2 boots the production configuration; handoff failure vllm-project#5 was MTP=3-specific, not the PCIe all-reduce. The traceback frame that pointed at custom_all_reduce was a graph-capture sync point, and dmesg carries the matching Xid 31 MMU faults on that run's PP0 pids. It does not pass the definition of done. Above ~7.7k prompt tokens the output stays fluent but misreads its own context; V1 is 3/3 clean on the identical probe, so this is a V2 regression rather than the driver, the GPU 6 Xid suspicion, or the 08-06/07 production corruption. VLLM_DCP_GLOBAL_TOPK=1 is recorded as a negative result. Also documents failure vllm-project#4 as a weighted-DCP-unaware block-size bug with a fix in the worktree, not a V2 limitation. Co-authored-by: Claude Opus 5 Signed-off-by: remichu <trantrungduc91@gmail.com>

…s to fp8-lane deepseek_v4 parser' (vllm-project#5) from fix/ctx645-dsv4-parser-backport into build/vllm-0.25.1-dsv4-sm121-stock

WoosukKwon added 4 commits March 3, 2023 04:16

Fix a bug in 1D shape

e5a1fa8

Minor

342275f

Minor

b91a2fa

Test iteration-level scheduling

4db2916

WoosukKwon merged commit 04e5acc into main Mar 6, 2023

WoosukKwon deleted the bugfix branch March 6, 2023 18:05

TheBloke mentioned this pull request Jul 20, 2023

Can't launch OpenAI API server on newly installed vLLM in Docker - fastchat not found #537

Closed

v1nc3nt27 pushed a commit to v1nc3nt27/vllm that referenced this pull request Sep 12, 2023

Merge pull request vllm-project#5 from ri938/more_improvements_awq

73db30f

More improvements awq

shanshanpt mentioned this pull request Nov 17, 2023

Run long conetxt error : CUDA error: an illegal memory access was encountered #1700

Closed

junior-zsy mentioned this pull request Nov 20, 2023

Error with 32k Long Text in chatglm2-6b-32k Model #1725

Closed

hongxiayang pushed a commit to hongxiayang/vllm that referenced this pull request Feb 13, 2024

Fix a bug in 1D input shape (vllm-project#5)

c639d4c

luo-cheng2021 pushed a commit to luo-cheng2021/vllm that referenced this pull request Mar 14, 2024

Merge pull request vllm-project#5 from slyalin/fixed_parameter_types

30605c8

Align optimum-intel based model signature with vLLM signature

luo-cheng2021 pushed a commit to luo-cheng2021/vllm that referenced this pull request Mar 25, 2024

Merge pull request vllm-project#5 from ilya-lavrenov/update-intel-opt…

8f016e0

…imum Install optimum-intel from latest main

dlopes78 mentioned this pull request May 8, 2024

[Bug]: VLLM + tritonserver #4695

Closed

Starmys pushed a commit to Starmys/vllm that referenced this pull request May 20, 2024

Merge pull request vllm-project#5 from wenxcs/fit-cluster-tests

de23377

Faster v2 hopper fused moe kernel configs

yuhuixu1993 mentioned this pull request Jun 2, 2024

[Bug]: loading squeezellm model #5190

Closed

oliver-li mentioned this pull request Jul 5, 2024

[Bug]: NCCL hangs and causes timeout #5484

Closed

This was referenced Jul 5, 2024

Support W4A8 quantization for vllm #5218

Merged

[Bug]: call for stack trace for "Watchdog caught collective operation timeout" #6042

Closed

ehuaa mentioned this pull request Jul 19, 2024

[Bug]: The vllm is disconnected after running for some time #5084

Closed

xinzaifeixiang1992 mentioned this pull request Jul 24, 2024

[Bug]: vllm-0.5.3.post1部署Qwen2-72b-instruct-awq模型，刚开始服务正常，但是并发高的时候就报错 #6734

Closed

alixiaodi mentioned this pull request Aug 2, 2024

[Bug]: #7072

Closed

Minami-su mentioned this pull request Aug 11, 2024

[Bug]: vllm is crashed on v0.5.3.post1 #7161

Closed

wangwensuo mentioned this pull request Aug 22, 2024

[Bug]: llama3-405b-fp8 NCCL communication #7775

Closed

zeroorhero pushed a commit to zeroorhero/vllm that referenced this pull request Sep 23, 2024

Merge pull request vllm-project#5 from KuntaiDu/kuntai-disagg-refactor

4db6446

Kuntai disagg refactor

SongXiaoMao mentioned this pull request May 13, 2026

[Bug]: MTP speculative decoding crash with illegal memory access on long sequences (Qwen3.6-27B-FP8, v0.19.1) #40756

Open

1 task

ibondarenko1 mentioned this pull request May 26, 2026

fix(processor): route MiMo-V2-Omni media fetch through MediaConnector #43117

Merged

arnavnagzirkar mentioned this pull request Jun 9, 2026

fix: [Feature]: CustomOp cleanup #44962

Closed

russellb mentioned this pull request Jun 10, 2026

Handle structured output grammar compilation failures #45012

Open

eastwood-c mentioned this pull request Jun 30, 2026

[Spec][V2] Support MTP speculative decoding under pipeline parallelism #46994

Open

connorcarpenter15 mentioned this pull request Jul 8, 2026

[Rust Frontend] Extend the existing gRPC service protocol #48033

Closed

5 tasks

dajinge mentioned this pull request Jul 9, 2026

[RFC]: Add API to restore free_block_queue allocation order for long-running RLHF / rollout sessions #40244

Open

1 task

JohnQinAMD mentioned this pull request Jul 29, 2026

perf(rocm): Kimi-K3 B1 route/AttnRes/tail/KDA fusion (vLLM stack) JohnQinAMD/vllm-amd#3

Closed

slokesha mentioned this pull request Aug 9, 2026

[XPU] Enable MiniMax-M3 on non-CUDA backends: portable fallbacks for CUDA-only paths #51578

Open

liranschour mentioned this pull request Aug 20, 2026

[Bug]: PD P2P supply and demand are computed independently, and nothing checks they match #53083

Closed

1 task

Ddhjx-code mentioned this pull request Aug 22, 2026

[Bug]: Scheduler permanently stops admitting requests once running + skipped_waiting ("deferred") reaches max_num_seqs - engine reports healthy, KV cache mostly free, only a restart recovers #53130

Open

1 task

Navigation Menu

Uh oh!

Uh oh!

Fix a bug in 1D input shape - #5

Fix a bug in 1D input shape#5
WoosukKwon merged 4 commits into
mainfrom
bugfix

WoosukKwon commented Mar 3, 2023

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Uh oh!

Uh oh!

Conversation

WoosukKwon commented Mar 3, 2023

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant