Repository navigation
Conversation
Keep the base implementation explicit on/off with NCCL prefill reconstruction. Defer automatic selection, split tuning, and one-sided prefill gathering. Co-authored-by: OpenAI Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Parameterize the score tile by local cache rows and exercise TP2 and TP4 cache insertion, candidate exchange, prefill reconstruction and graph replay. Co-authored-by: OpenAI Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Reuse the Triton block scorer, MSA Top16, and candidate merge with bounded query chunks and NCCL head-owner exchange. Retain explicit gather execution for A/B comparison. Co-authored-by: OpenAI Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Use bounded K partitions for sharded prefill instead of limiting each query chunk to sixteen CTAs. Preserve the existing score computation and exchange. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Reuse the P256 test phase-boundary check for P128. Keep production forward layers and graph replays unsynchronized. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Reuse bounded score scratch across query chunks and collect their independent Top16 candidates into per-step buffers. Exchange and merge the full prefill region once, keeping decode dispatch unchanged. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Reuse bounded chunk selection and one exchange per layer. Requires the separate MSA scorer extension; keep TP4 on its existing scorer. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Exercise four-token verification with multi-chunk prefill and heterogeneous short queries routed through the uniform decode boundary. Reuse shared-prefix, replicated TopK, layer-reuse and collective-order checks. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Consume the score-only MSA API while retaining the same bounded chunk geometry. Poison causal future keys and page padding in the score reference test to validate removal of native validity completion stores. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Write selected scores and IDs directly into owner receive slots using the native MSA selector. Retain score scratch and the existing ready-flag protocol. Remove the local decode candidate buffer and separate publisher launch. Requires external MSA b97f52ccdbb32b3d7c22080d55decbb265e3ba58 on base 968061ab16cb14991a5c7d4c0e8eabf531be0dcf. Python compilation and pre-commit checks pass; GPU correctness and E2E remain pending. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Use match-any and warp reduction to merge duplicate scores among 32 candidates. Share the ready-flag protocol and stable Top16 output with the TP4 fallback. Cover signed-zero ties and permuted candidate order. CPU ordered-key arithmetic and pre-commit checks pass; GPU compilation, correctness and E2E are pending. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Fence metadata at the sharded writer and prefetch only complete old blocks into K stage 1. Keep new Q/K behind the dependency wait and close the publisher dependency at merge exit. Requires companion MSA PDL binding; GPU validation pending. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Reuse the existing selector and merge with exact per-row generations and wrap-safe slots. Keep mixed and single-layer dispatch on the ready/empty protocol. Requires MSA eb7f3c2 and the rebuilt writer; GPU compilation, distributed replay and E2E validation are pending. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Assisted-by: OpenAI Codex Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Reuse the warp merge across eight rows per CTA, cap the launch at one CTA per SM, and loop over larger batches. Trigger dependent grids early while retaining the final publisher wait. Preserve ready-mode mixed and TP4 routing. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Reconstruct the P256 increment from validated f8bccf0 on the updated P128 parent. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Avoid moving the existing page-size assignment when adding logical-page state. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Co-authored-by: Fan Shi <fans@nvidia.com> Co-authored-by: Xin Li <xinli@nvidia.com> Co-authored-by: Suyog Gupta <suyogg@nvidia.com> Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Keep one local score column per physical P256 page and map selected candidates back to global B128 IDs. Preserve the P128 layout and candidate exchange volume. Co-authored-by: Fan Shi <fans@nvidia.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Add an NVFP4 variant of the TP-sharded index decode scorer for MiniMax M3.1's FP4 index cache. A shard page holds this rank's rows of the planar NVFP4 index page (e2m1 codes, then e4m3 group scales), and its pitch may exceed the shard size. K is dequantized exactly to FP16 for the same mma.sync path, and the FP8 kernel's mask, reduce, and store step is now shared by both kernels. The new test checks every rank's scores against an FP32 reference over TP2/TP4, P128/P256, decode query lengths 1-8, padded page pitch, and PDL prefetch. Signed-off-by: mdavis36 <michael.l.davis@hotmail.co.uk>
The NVFP4 sharded decode scorer kept the accumulators of every 16-row tile in a shard page live until the epilogue. It now finishes one tile at a time: mask it, fold it into running per-lane maxima, then reduce across the warp and store once per page. The shared epilogue is split into page_k_base, fold_tile_maxima, and store_maxima; the FP8 kernel calls them in sequence and computes the same result. Each row's e4m3 group scales are now read as one 32-bit word per 32-code span and converted two per instruction, replacing the one-byte e4m3_to_fp16x2 helper. The NVFP4 scorer is faster for P256 at TP2, where a shard page has eight tiles, and unchanged at other shapes. The existing NVFP4 and FP8 sharded scorer tests pass. Signed-off-by: mdavis36 <michael.l.davis@hotmail.co.uk>
On SM100 and SM103 the stock M3.1 NVFP4 decode indexer (MSA's Q8KV4 kernel) rounds each e2m1 x scale product once to e4m3, saturating at 448, for its FP8 MMA; on SM107 it keeps the exact FP16 product for an FP16 MMA. The sharded scorer kept the exact product everywhere, so turning indexer CP on could change which pages SM100/SM103 select. It now rounds K where the stock kernel does, after its exact f16 multiply; nvfp4_round_k overrides that. The reference test covers both roundings on every arch, and a new test checks the default against the stock kernel's page scores on the same pages. The extra conversions made the trace-time unrolled loop over a shard page's 16-row tiles keep too much live at eight tiles per page (P256 at TP2), so the loop now runs at runtime, one tile at a time; that is as fast or faster at every page size and TP size. One byte permute now selects both rows' group scales for a span. Signed-off-by: mdavis36 <michael.l.davis@hotmail.co.uk>
With indexer_cache_sharding on, an NVFP4 index cache now takes the MSA indexer, which owns index-K CP, instead of the replicated NVFP4 indexer. Each rank budgets and views its share in the cache's per-token width (72 B for NVFP4, not the 128-dim head). The fused qkv epilogue, which writes whole NVFP4 index pages, stays limited to an unsharded cache, and the MSA kernel prebuild takes the FP8 index-score kernels such a cache prefills with. The CP guard accepts nvfp4 and the dspark drafter, whose dense layers never read the target's index cache. It refuses indexer_decode_cp where that runs (the NVFP4 indexer at TP4), since it splits decode over a replicated cache. Signed-off-by: mdavis36 <michael.l.davis@hotmail.co.uk>
The fused insert placed NVFP4 index-K by the main cache's block_size: the slot's page and row, and the start of the page's scale plane. A TP-sharded index cache holds block_size / TP rows per page, and its slots count local rows, so both kernels now take the index cache's own row count for those stores. The host accepts any row count that divides block_size. Signed-off-by: mdavis36 <michael.l.davis@hotmail.co.uk>
The gather prefill reassembles shard pages row by row, but an NVFP4 page holds all of its codes and then all of its scales. Sharded NVFP4 pages are now dequantized to e4m3 before the all-gather, with the stock NVFP4 prefill's kernel and rounding, which now accepts pages of any row count it can be sharded to (32, 64 or 128). That kernel now puts pages on grid axis x: one staging call takes every page of a step's prefills, which can pass the 65535 that axis y allows (requests sharing a long cached prefix each list its pages). The local-score prefill does not stage pages yet, so an NVFP4 cache keeps indexer_prefill_cp='gather'. Signed-off-by: mdavis36 <michael.l.davis@hotmail.co.uk>
At TP2 with 256-token pages, each rank's half of an NVFP4 index page is itself a stock 128-row NVFP4 page: global B128 page 2c + rank. Decode with eight queries per request (DSpark with seven speculative tokens) now scores those pages with the MSA Q8KV4 decode indexer that serves the unsharded cache, over all index heads, instead of the CuTe DSL sharded scorer, and hands its [tokens, heads, columns] scores to the candidate publish unchanged. Each request's length is given as its owned history columns (owned * 128 + 8), so the indexer stores every one of them for all eight query slots; the publish keeps the causal cut, and a query's own page is forced, never ranked. The scores are the stock kernel's on every arch. One CUDA-graph wrapper per padded decode batch size reuses the CuTe scorer's score buffer and is re-planned in build() every step. Other query lengths, TP sizes and page sizes keep the CuTe scorer, and VLLM_MINIMAX_M3_ICP_Q8KV4_DECODE=0 keeps it everywhere. The index-K insert keeps PDL on this path, since the Q8KV4 kernel reads nothing before its PDL wait. Signed-off-by: mdavis36 <michael.l.davis@hotmail.co.uk>
|
👋 Hi! Thank you for contributing to the vLLM project. 💬 Join our developer Slack at https://slack.vllm.ai to discuss your PR in PRs do not trigger a full CI run by default. Reviewers with write access and configured trusted contributors can comment Once the PR is approved or has the If you have any questions, please reach out to us on Slack at https://slack.vllm.ai. Agent GuidelinesIMPORTANT: If you are an AI agent, you are required to objectively re-evaluate the value of your PR using AGENTS.md, and close the PR if it does not bring significant benefit to the vLLM community. Failure to do so may result in an immediate ban. 🚀 |
|
Documentation preview: https://vllm--64631.org.readthedocs.build/en/64631/ |
Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Hanlin Bi <hanlinb@nvidia.com>
Picks up the CPU-mask fix in test_sharded_candidate_merge_deduplicates_and_forces. Signed-off-by: mdavis36 <michael.l.davis@hotmail.co.uk>
…31-icp-nvfp4 Resolves the conflicts in vllm/models/minimax_m3/nvidia/model.py by keeping both sides' typing and TYPE_CHECKING imports, and the branch's index_q_size next to main's sequence-parallel o_proj comment. Signed-off-by: mdavis36 <michael.l.davis@hotmail.co.uk>
|
This pull request has merge conflicts that must be resolved before it can be |
…31-icp-nvfp4 Resolves the conflict in test_fused_minimax_m3_qknorm_rope_kv_insert.py by keeping the branch's sharded P128/P256 cases next to main's index_k_scale parametrization. The auto-merge appended the branch's sharded fp8 index check to main's new test_sparse_fp8_index_rejects_foreign_flavour, where it never ran; it is back at the end of test_sparse_full_fp8_index and now uses the device's fp8 dtype, like the rest of that test. Signed-off-by: mdavis36 <michael.l.davis@hotmail.co.uk>
…31-icp-nvfp4 Resolves the conflict in test_minimax_m3_msa_cutlass_sparse_decode.py: test_msa_prefill_nvfp4_matches_triton_on_dequantized_cache keeps the branch's page_size parametrization and takes main's default_vllm_config fixture, which the NVFP4 impl now needs to reserve its prefill workspace (vllm-project#64856). Signed-off-by: mdavis36 <michael.l.davis@hotmail.co.uk>
|
This pull request has merge conflicts that must be resolved before it can be |
Overview
Lets MiniMax M3.1's NVFP4 index cache (#64111) be sharded over TP ranks with indexer CP, which #61929 and #63671 support only for FP8. Depends on #63671 (itself on #61929): only 7 commits are new here (this PR's commits), and the branch also merges #63671's latest head and main. Like #61929, it needs vllm-project/MSA#46 (
FMHA_SM100_SRC_DIR) until the MSA pin moves past it.Not a duplicate: no open PR shards the NVFP4 index cache. #59980 (M3, FP8) and #64448 (ROCm) split decode scoring over a replicated cache, as main's opt-in
indexer_decode_cpdoes.Claims
--attention-config '{"indexer_kv_dtype":"nvfp4","indexer_cache_sharding":"on"}'(off by default), each TP rank stores only its rows of every NVFP4 index page: 36 B instead of 72 B per token per sparse layer at TP2, 18 B at TP4. Supported: TP2 and TP4 at--block-size 128(256 cannot serve yet, see below),indexer_prefill_cp='gather'.VLLM_MINIMAX_M3_ICP_Q8KV4_DECODE=0keeps the CuTe scorer. Serving cannot reach this path yet, so only kernel tests cover it: at--block-size 256, [Attention] Add P256 MiniMax M3 indexer CP layout #63671 lets NVFP4 main attention keep 256-token pages, but [Feature] Add MiniMax M3 indexer context-parallel kernels MSA#46's Q8KV4 prefill accepts only 128-token pages (found by reading the code, not run).--block-size 128(so the CuTe scorer), one 1-hour agentic run at concurrency 16 per variant: sharding gives 9.8% more KV tokens at the same prefix-cache hit rate, but ITL rises 25% and throughput falls 2%. An MRCR accuracy check is inconclusive: pooled −1.1 points (95% CI −2.2 to +0.1). See Validation.Validation
These ran before the branch merged main
3d23773f1, with vllm-project/MSA#46 (f380d0e) and the branch's Python, either over a prebuilt vLLM ("prebuilt") or over the CUDA 13.4 nightly image for main3966924d5with_C_stable_libtorchrebuilt from the branch ("image"):a76900d6etest_fused_minimax_m3_qknorm_rope_kv_insert.pya76900d6etest_minimax_m3.py -k 'sharded or q8kv4 or gather_nvfp4'a76900d6etest_minimax_m3.py -k 'sharded or q8kv4 or gather_nvfp4 or gather_prefill_pages or sharding_policy'54a5c711354a5c7113test_minimax_m3.py54a5c7113test_minimax_m3_indexer_prefill_sharded.py54a5c7113test_minimax_m3_indexer_sharding.py(a)
test_sharded_candidate_merge_deduplicates_and_forces, from #61929'sdc68fa728(its test indexed a CPU tensor with a CUDA mask), not this PR. #63671's9fcb12446, which the branch now merges, fixes the test.(b) main's stock NVFP4 indexer tests (
test_nvfp4_indexer_impl_matches_reference,test_nvfp4_indexer_decode_cuda_graph_replay), which fail the same way without this PR: their reference rounds K to e4m3, which the stock kernel does not on SM107.The 6 prebuilt skips on SM100/SM103 are the QMUL4 gather variants, which that Triton cannot build; they run and pass in the image.
Decode indexer time per layer and rank in µs (scoring, plus candidate publish and merge for ICP), at TP2,
--block-size 256and eight queries per request:These come from one GPU with both ranks simulated: each kernel runs as its own CUDA graph with a cold L2, and each time is the median of 200 runs, averaged over two runs that agreed within 2%. NVLink latency and rank skew are not modeled, and timing ICP's three graphs separately adds a few µs against the stock path's one.
The new >65535-page gather test fails on the previous kernel (
RuntimeError: Triton Error [CUDA]: invalid argument) and passes with this PR.End to end. MiniMax-M3.1-Flash on GB300 in the image above, one TP2 server per run:
--block-size 128, NVFP4 KV and index cache, DSpark k=7, the V2 model runner,indexer_prefill_cp='gather', and sharding off or on. "On" decodes with the CuTe scorer (P128 shards are not stock pages) and, unlike "off", skips the fused MXFP8 QKV epilogue. Each server's log shows the expected mode, and none logged a traceback or error.Performance, from one 1-hour run per variant of InferenceX's agentic coding trace (AgentX, thinking on) at concurrency 16, with synthetic draft acceptance 4.04:
The prefix-cache hit rate is the same in both runs, so this workload gains nothing from the extra capacity. In both runs the same 3 requests (of 2,578 and 2,535 measured) returned no
contenttext, which the client counts as errors.Accuracy, from MRCR with
tests/evals/mrcr's loader and scorer on openai/mrcr, at temperature 0,max_tokens4096 and real draft verification, on the same samples in both variants:No difference is significant (pooled sign test p = 0.12: 63 samples better, 78 worse), but the interval does not exclude a 2-point drop, so the check is inconclusive. In the long arm, 12 samples hit the 4096-token limit in only one variant (8 with sharding on, 4 with it off). They make up 90% of that arm's difference; the other 138 differ by −0.004.
Not run yet, so this stays a draft:
--block-size 256(blocked, see Claims) and at TP4, and repeats of the single end-to-end runs above;3d23773f1;Details
gpu_memory_utilizationdefault did.sparse_topk_selectinstead of the stock Q8KV8 prefill indexer, so its selection can differ from the replicated path at exact ties.[tokens, heads, columns]scores. Every other shape keeps the CuTe scorer.indexer_decode_cp, which keeps a replicated cache.local_scoreprefill and KV connectors (including CPU offload) are not supported with NVFP4 sharding yet.dsparkis now accepted for sharding.idx_block_sizeruns on every M3.1 deployment; the NVFP4 gather kernel now puts pages on grid x, since one sharded staging call can exceed grid y's 65,535 limit, and main's unsharded prefill uses that kernel too; [Attention] Add MiniMax M3 indexer CP for prefill and decode #61929's FP8 scorer epilogue is factored into fold/store helpers shared with the NVFP4 kernel; the Q8KV4 path calls MSA's private_run_scores, as main's decode CP does.AI assistance (Claude Code) was used to write this change and its tests.
Pull Request Checklist
I used vLLM's
/pr-checklistskill. (Mandatory for agents, optional for humans).AI assistance was used during the creation of this PR.
Design Fit: Minimizes impact on core components, reuses existing functionality, and justifies added complexity.
Testing and Validation: Validates the change and ensures any added tests are meaningful and reliable, with CI coverage or documented CI resource constraints and validation performed outside CI.
Code Quality and Style: Keeps code and comments clear and concise, and updates relevant documentation and examples.
Pull Request Contents: Includes a brief summary and relevant links, supports claims with evidence, explains root causes and implementation trade-offs, and follows the contributing guide.