[Bugfix][DSv4] SM12x FlashInfer + DSpark capture for GB10 - #1
Merged
Merged
Conversation
Stock v0.28 already selects FlashInfer sparse MLA on SM12x, but two gaps block DeepSeek-V4-Flash-0731 on dual DGX Spark: - compute_fp8_einsum_recipe treated major>=10 as SM100, packing UE8M0 scales that become garbage on GB10. Use the Hopper FP32 recipe. - mhc_pre_broadcast_tilelang always called DeepGEMM prenorm. Fall back to the existing torch prenorm GEMM when DeepGEMM is unavailable. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
FlashInfer 0.6.17 ships SM12x DSV4 decode cubins for top_k in
{128, 512, 1024}, not DSpark k=5's aligned width 192. Snap the
non-causal SWA index list up to the next cubin so init and the
kernel see the same supported shape; the logical window stays 128.
Also reshape spec-decode queries to [batch, next_n, heads, 512] and
route <=64-token prefill segments through decode-form calls so
FlashInfer does not misroute DSpark drafts to the SM120 prefill
kernel (vllm-project#52499).
Co-authored-by: Cursor Grok 4.6
Signed-off-by: Cursor Agent <cursoragent@cursor.com>
cursor
Bot
changed the title
Keep the v0.28 default at 0.6.17. Add a dated vllm-project#4380 nightly overlay so GB10 can install native DSV4 topk=192 without tracking latest. Direct GitHub URLs survive nightly-index rotation. No cu130 aarch64 jit-cache on this date; uninstall 0.6.17 jit-cache so JIT rebuilds. Replace with 0.6.18 stable when it ships. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Package the Spark 0.28 + DSpark + FlashInfer 0814 stack as an aarch64 image built on the first GB10, then docker save/load to the second node instead of rsyncing a venv. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Idle Spark compiles can use more ARM cores; 8 stays short of nproc (20) so nvcc is less likely to OOM unified memory. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Keep host/IB/NCCL/offline mounts from the 0.1.1 command. Drop B12X and nvfp4_ds_mla; stock 0.28 uses fp8_ds_mla and auto MoE. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
build.sh from-scratch rebuilds torch + vLLM. pack-venv.sh copies ~/.venv/vllm028 and relocates shebangs/.pth. Optional INSTALL_FLASHINFER_NIGHTLY=1 only overlays the 0814 wheels. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Host venvs are editable. PEP 660 finders keep the /home path after COPY, so /opt/venv/bin/vllm raises ModuleNotFoundError. Rewrite finder .py files, drop leftover vllm editables, write a relocated .pth, and repair at container start so the current image can launch without a from-scratch rebuild. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Packed image had libibverbs1 but no /etc/libibverbs.d providers, so NCCL_NET=IB failed at ncclCommInitRank. Add mlx5/rdma packages and bind-mount the host verbs plugins for the already-built image. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Apt libibverbs1 is not enough for NCCL_NET=IB. collect-ibverbs.sh copies this node's /etc/libibverbs.d and mlx5 libs; pack-ibverbs.sh injects them into the existing image without recopying the venv. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Default capture stride is 8, so --max-cudagraph-capture-size 36 was truncated to 32. List 36 explicitly for DSpark 6*(1+5). docker run -d so head/worker stay up without a tty. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Profile dummy-run called cutlass_scaled_mm; dispatch_scaled_mm requires float scales, but DeepSeek-V4 Flash stores UE8M0. SM12x was treated as SM100-capable Cutlass. Disable that kernel on SM12x and upcast UE8M0 scales to fp32 for Triton/Cutlass. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Install the stock b12x extra into vllm-gb10:dspark without recompiling vLLM, and launch with --linear-backend b12x --moe-backend b12x plus a persistent CuTe JIT cache. Not Anemll flashinfer_b12x. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Compiled DeepGEMM einsum asserts m/n/k on GB10. Keep the Hopper (1, 128, 128) recipe and dispatch o_proj to the vllm-project#52357 Triton fallback. Do not Marlin/Humming-pack is_bmm wo_a weights. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Python nvidia-cutlass-dsl emits enable-pyir=false that the installed cute-to-nvvm pass rejects on sm_121a. Gate all in-tree DSv4 CuteDSL compiles (indexer Q, compressor, K-gather) and keep the Triton path. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
PIECEWISE CUDA-graph capture used min(num_tokens, max_num_seqs) dummy requests, so size 32 with --max-num-seqs 6 became a ragged 32-token / 6-request decode batch. SM12x DSpark decode requires uniform next_n and asserted. Build a uniform dummy (largest divisor of the capture size) and fall back to per-request [1, q_len, ...] decode-form when a batch is still ragged. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
cursor
Bot
force-pushed
the
cursor/dsv4-flash-sm12x-spark-df88
branch
from
August 28, 2026 09:15
7947f59 to
544e72a
Compare
Packed GB10 venvs are uv-managed and have no pip module, so pack-b12x.sh failed with "No module named pip". Bootstrap uv and install official b12x==1.2.6 with the current torch pin. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
CUDA-graph capture succeeded, then kernel warmup compiled the ll_bf16 router GEMM via cute.compile. SM12x still hits the cute-to-nvvm enable-pyir ICE. Gate all in-tree CuteDSL JIT (including ll_bf16, FA4 warmup, and generic CuteDSL warmup) behind cutedsl_jit_supported(); GateLinear keeps cuBLAS/F.linear. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
b12x==1.2.6 lazy-loads attributes and has no __version__, so pack-b12x.sh failed after a successful uv install. Read the dist version from importlib.metadata. Official extra may pin nvidia-cutlass-dsl==4.6.2 over a packed 4.7.0. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
DeepSeek-V4-Flash uses HF config.json. The tokenizer mode is deepseek_v4, but that name is not a config parser, so headless workers failed ModelConfig validation. Map the alias to HFConfigParser. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Official b12x==1.2.6 compiles sm_120 TMA dense GEMMs that IMA on GB10 (sm_121a) during PIECEWISE capture of shared-expert gate_up. Keep --moe-backend b12x; --linear-backend b12x falls back to Triton for block-scaled FP8. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
After skipping b12x block-FP8 on SM121, kernel selection picked Marlin weight-only FP8. PIECEWISE dummy capture then IMA'd in shared-expert gate_up on aux_stream (aten::empty in marlin_gemm). SM12x has native FP8; use Triton. Keep shared experts on the main stream. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
PIECEWISE capture of size 8 is 4 reqs × 2 tokens. FlashInfer SM120
batched decode IMA'd on next_n=2 after next_n in {4,6,8} succeeded.
Use per-request [1, q_len] on SM12x unless next_n is DSpark k+1.
Also disable MHC TileLang PDL and the n_splits=4/8 small-FMA path
on SM12x (n_splits=1 survived the larger captures).
Co-authored-by: Cursor Grok 4.6
Signed-off-by: Cursor Agent <cursoragent@cursor.com>
PIECEWISE 50% is size 8. After skipping batched next_n=2, capture
still IMA'd: dummy stayed 4 reqs × 2 tokens, and MHC still saw 8
tokens (size 16 succeeded in the same compiled kernel).
- Dummy packing avoids tokens_per_req in {2,3} (8 → 2×4)
- SM12x FlashInfer splits leftover q_len=2 into two q_len=1 launches
- SM12x MHC pads token dim <16 to 16
- Disable DeepGEMM PDL on SM12x
Co-authored-by: Cursor Grok 4.6
Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Size-8 PIECEWISE warmup still IMA'd after dummy packing became 2×4 and MHC padded to 16. The stack is DeepseekV4 FFN -> B12xExperts -> TP all-reduce, so the faulting kernel is 8-token MXFP4 MoE (size 16 already succeeded). Pad hidden/topk to 16 on SM12x inside FFN and B12xExperts, then slice back. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Padding hidden_states 8→16 made topk_softplus_sqrt see 16 tokens while ForwardContext.is_padding was still length 8 (is_padding size mismatch, expected: 16). Extra aligned rows are marked True and the context mask is restored after experts(). Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Mixed-warmup's 2-token seed prefill is routed through decode-form FlashInfer (SM120 prefill requires >64 tokens). Splitting q_len=2 into two q_len=1 launches IMAs because the second token's KV is not written yet. Pad to a capture-proven width and seed mixed warmup at q_len=4. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
The mixed step was 1 decode + 15 prefill. SM120 decode-form IMA'd below 16; snap the prefill width and keep scheduled tokens consistent. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
DSpark sample_from_anchor uses q_len=5. The first FULL dummy is then 25=5x5, a shape main capture never ran; 16:26 died at 0/4 before FlashInfer could pad. Align capture to decode q_len=6 so the first dummy is the already-green 36=6x6. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
cursor
Bot
changed the title
Aligning DSpark capture to q_len=6 would still stage 18=3x6 and 12=2x6, token counts main capture never ran. Keep only 36, 24, and 6 (MHC pads 6 to 16) so capture can finish and mixed warmup can pad q_len=2 to one [1, 6] launch. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
16:26 died at DSpark 0/4 with no traceback. Synchronize before the first dummy so a pending main-capture IMA is not blamed on DSpark, and after each dummy so the failing tokens/reqs/q_len line is the one that actually IMA'd. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
SM12x treat_short_extends=False makes split_decodes_and_prefills require CommonAttentionMetadata.is_prefilling. DSpark capture built metadata without it, so the first dummy AssertionError'd at 0/5 (16:41). Pass make_dummy's all-False flag, matching the target-model capture path. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Capture dummies were the 16:41 AssertionError. Runtime _build_draft_attn_metadata had the same hole and would fire on the first propose after capture, including mixed warmup if it runs the speculator. Draft steps are decodes. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
treat_short_extends=False still uses the flag to keep mixed-warmup q_len=2 on the prefill pad-to-[1, 6] path. A missing flag must not AssertionError the worker (DSpark capture 16:41). Treat None as no extra prefills. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
C128A used the default treat_short_extends=True. With DSpark's threshold that swallows a 2-token mixed-warmup seed as all-decode, while SWA keeps it as prefill. FlashInfer then asserts c128a_prefill_topk_indices is None before it can pad q_len=2 to [1, 6]. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
synchronize() after each dummy was meant to surface a pending IMA. Inside torch.cuda.graph it raises cudaErrorStreamCaptureUnsupported (Spark 16:53, DSpark 0/3). Sync only after the eager warmup replay. Keep the pre-dummy sync before capture starts. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Spark 17:03 confirmed DSpark 3/3 and padding prefill q_len=2 -> 6 query.shape=(1, 6, 32, 512). That decode-form first-prefill then IMA'd (reported at C4A compute_global_topk). DSpark [1, 6] decode dummies were fine, so first prefills now pad to 65 and use the SM120 prefill kernel. C4A topk also clamps req/block indices. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
The 65-token SM120 prefill-kernel workaround redefined the requested launch. Mixed-warmup first prefills again use one decode-form [1, 6] (never 2→4). C4A block-table clamp stays. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Spark seed is C4A (compress_ratio=4). Guard that path as one decode-form [1, 6] launch, never [1, 4]. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Spark 17:58 confirmed padding prefill q_len=2 -> 6 query.shape=(1, 6, 32, 512), then IMA at MoE all_reduce. Decode-form q_len=6 reads six KV / C4A rows; insert only wrote two. Extend same-block slots and repeat the last row so the launch stays one [1, 6] (never 2→4). Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Spark 18:13: PIECEWISE 5/8 died in compressor sm12x_extend_prefill_slots via bool(torch.all) on a tokens=4 decode dummy. Fill only eager all-prefill; extend slots with torch.where (no GPU→CPU sync). Keep mixed-warmup q_len=2 -> [1, 6]. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Spark 22:57: capture 8/8 + 2/2 + DSpark 3/3, insert filled 6 SWA slots, FlashInfer launched [1, 6], then IMA at MoE all_reduce. Padded first-prefill still has a 2-token indexer topk; repeating those extra-sparse rows into a live C4A cache is the remaining hole. Keep the SWA [1, 6] pad; detach extra cache only for padded prefills. DSpark decode 5->6 keeps C4A. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Spark 23:42: 5878020 captured 8/8 + 2/2 + DSpark 3/3, launched SWA-only [1, 6] x3, then IMA at compute_global_topk Triton on the next layer. Compressor still wrote 6 C4A tokens and prefill still built unused extra-sparse metadata from a 2-token indexer. Keep SWA insert and the [1, 6] pad; do not pad C4A writes or run compute_global_topk when every prefill will drop extra cache. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
info_once cannot take a list; Spark unit tests hit TypeError before the SWA-only [1, 6] launch. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Spark 23:54: 61ccf6a skipped C4A and launched SWA-only [1, 6] through many layers, then IMA at wo_b all_reduce. MHC/MoE already pad token dim <16; o_proj still ran inverse-RoPE + FP8 einsum on 2 tokens. Pad o and positions to 16 and slice the output back. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Spark 00:05: ee2ff4e padded o_proj to 16, capture stayed green, then the same mixed-warmup seed IMA'd after many SWA-only [1, 6] launches. Padding query rows was not enough: each real SWA row already has 126-127 -1 tails (swa_len 1-2, width 128). FlashInfer SM120 decode cubins gather the full width first. Repeat the last valid slot per row; keep swa_topk_lens as the softmax mask. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Ruff E501 / format on the sentinel replacement and nearby lines. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Same serve as docker/gb10/run.sh plus CUDA core dump env, a host dump mount, optional CUDA_LAUNCH_BLOCKING, and optional compute-sanitizer. Daily serve stays on run.sh. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
cuda-gdb is /usr/local/cuda/bin/cuda-gdb inside vllm-gb10:dspark, not on the Spark host PATH. The serve container dies after IMA, so open the host dump with a one-shot container. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Spark 01:12 coredump: WARP_ILLEGAL_ADDRESS in b12xmoe MoEDynamicKernelSilu. Head 07:24 still IMA'd at MoE all_reduce after SWA-only [1, 6]. VLLM_MOE_SKIP_PADDING writes topk_ids=-1 on SM12x alignment rows. Capture dummies mark every row padding so the kernel no-ops; mixed-warmup has 2 real rows plus 14 -1s. Clamp sentinels to expert 0 and zero those weights. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Cursor Agent <cursoragent@cursor.com>
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
This was referenced Aug 31, 2026
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
2 participants
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
Make stock vLLM 0.28+ able to serve official DeepSeek-V4-Flash-0731 on 2× DGX Spark (GB10 / SM121) with DSpark (
k=5).AI assistance: Cursor Grok 4.6 was used. A human submitter must review every line before merge.
Spark evidence
01:12
debug.shcoredump:07:24 head
run.sh: capture green, mixed-warmup seed[1, 6], IMA reported at MoEall_reduce(NCCL reporter). Same kernel.VLLM_MOE_SKIP_PADDINGwritestopk_ids=-1on SM12x alignment rows. Capture dummies mark every row padding → all-1→ SiLU no-ops. Mixed-warmup has 2 real rows + 14-1s → gather IMA.Fix (SHA
342375171)sm12x_replace_moe_topk_sentinelsclamps-1to expert 0 and zeros those weights beforeB12xExperts.apply. Daily serve stays onrun.sh(notdebug.sh).Test Result
Spark: pull
342375171, rsync, restart worker then head withrun.sh.