[Kernel][Perf] Add fused CUDA post-conv MTP decode kernel for Qwen3.5 GDN by Jie-Fang · Pull Request #51674 · vllm-project/vllm

Jie-Fang · 2026-08-10T12:49:25Z

Purpose

Speed up Qwen3.5 (GDN linear attention) MTP speculative decode on Blackwell. During
MTP decode, the Triton path launches a chain of small kernels per step (gating,
delta-rule recurrence, state rewind/update, gated RMSNorm), which leaves the GPU
latency-bound at decode batch sizes.

This PR adds a single fused CUDA kernel, fused_gdn_decode_post_conv_mtp, that
consumes the post-convolution mixed_qkv and performs the entire GDN MTP decode
step in one launch:

SiLU activation on q/k/v and gating (softplus/sigmoid on a, b, A_log,
dt_bias)
the gated delta-rule recurrence over up to 8 speculative tokens per request,
with rewind to the last accepted token via num_accepted_tokens
recurrent state update in place (BF16 or FP32 state)
gated RMSNorm epilogue, with output precision matched to the Triton path

Test Plan

New/extended tests (all require SM100):

pytest tests/kernels/mamba/test_gdn_fused_mtp.py -v                                                                                                                                                                                                                                     
pytest tests/kernels/test_fused_gdn_post_conv.py -v                                                                                                                                                                                                                                     
pytest tests/v1/attention/test_gdn_metadata_builder.py -v                                                                                                                                                                                                                               
pytest tests/test_envs.py -k "gdn"                                                                                                                                                                                                                                                      
pytest tests/compile/test_config.py

test_gdn_fused_mtp.py covers the model-path dispatch for pure spec-decode,
mixed, prefill, and regular-decode batches against the Triton reference.
test_fused_gdn_post_conv.py::test_fused_gdn_decode_post_conv_mtp_ratio8 covers
the kernel directly for BF16/FP32 state, ragged acceptance patterns, and TP head
counts (tp4/tp16).

Test Result

On 1x B200 (SM100), CUDA 13.1, torch 2.13.0+cu130:

tests/kernels/mamba/test_gdn_fused_mtp.py: 6 passed
tests/kernels/test_fused_gdn_post_conv.py: 70 passed
tests/v1/attention/test_gdn_metadata_builder.py: 10 passed
tests/test_envs.py -k gdn: 1 passed
tests/compile/test_config.py: 44 passed

Microbenchmark

1x B200, CUDA 13.1, torch 2.13.0+cu130. Both variants measured over the
same kernel boundary: post-conv mixed_qkv in, gated+normed
core_attn_out out, recurrent state updated in place.

Triton path = fused_sigmoid_gating_delta_rule_update + gated RMSNorm
(rmsnorm_fn), exactly as dispatched by _forward_core_fused_norm today.
Fused path = single fused_gdn_decode_post_conv_mtp launch.
Timed by capturing 20 iterations in a CUDA graph and replaying 50x with CUDA
events (production MTP decode runs under full CUDA graphs, so graph-replay
GPU time is the relevant metric; CPU launch overhead excluded for both).
cold rotates input/state buffer sets so the working set exceeds 4x L2
(126 MB), i.e. state comes from HBM as in real serving where other layers
evict L2 between calls; warm reuses one buffer set (L2-hot upper bound).
K=V=128, 4 MTP tokens/request, num_accepted_tokens=2. Context length does
not affect either variant (fixed-size recurrent state), so results apply to
any seqlen (e.g. 8192).

Cold-cache speedup is 1.17x-2.20x, largest in the latency-sensitive
BS=4-32 decode regime.

TP=4 (H=4, HV=32), bfloat16 recurrent state

BS	Triton cold (us)	Fused cold (us)	Speedup	Triton warm (us)	Fused warm (us)	Speedup
1	20.31	12.19	1.67x	16.49	11.57	1.42x
4	27.91	12.99	2.15x	20.65	11.73	1.76x
8	35.69	17.89	2.00x	27.26	16.59	1.64x
16	46.66	32.32	1.44x	41.44	30.60	1.35x
32	71.85	51.75	1.39x	71.99	51.62	1.39x
64	120.65	97.12	1.24x	120.61	96.93	1.24x
128	221.13	188.16	1.18x	221.11	188.02	1.18x

TP=4 (H=4, HV=32), float32 recurrent state

BS	Triton cold (us)	Fused cold (us)	Speedup	Triton warm (us)	Fused warm (us)	Speedup
1	23.50	13.73	1.71x	17.74	12.70	1.40x
4	32.04	14.57	2.20x	24.12	12.82	1.88x
8	42.69	20.78	2.05x	33.78	18.68	1.81x
16	60.27	38.14	1.58x	60.01	37.70	1.59x
32	96.82	65.27	1.48x	97.08	65.15	1.49x
64	167.63	130.97	1.28x	167.27	126.25	1.32x
128	311.52	267.23	1.17x	314.21	266.75	1.18x

TP=16 (H=1, HV=8), bfloat16 recurrent state

BS	Triton cold (us)	Fused cold (us)	Speedup	Triton warm (us)	Fused warm (us)	Speedup
1	15.87	11.52	1.38x	15.92	11.57	1.38x
4	19.64	12.32	1.59x	16.46	11.57	1.42x
8	23.52	12.67	1.86x	17.92	11.61	1.54x
16	27.56	12.92	2.13x	20.64	11.69	1.77x
32	35.39	17.82	1.99x	27.16	16.60	1.64x
64	46.45	32.27	1.44x	40.72	30.57	1.33x
128	73.57	51.77	1.42x	73.69	51.61	1.43x

TP=16 (H=1, HV=8), float32 recurrent state

BS	Triton cold (us)	Fused cold (us)	Speedup	Triton warm (us)	Fused warm (us)	Speedup
1	17.44	12.59	1.38x	17.71	12.61	1.41x
4	23.30	13.71	1.70x	17.92	12.68	1.41x
8	26.18	13.93	1.88x	20.08	12.71	1.58x
16	31.77	14.50	2.19x	23.99	12.81	1.87x
32	42.81	20.84	2.05x	34.81	19.05	1.83x
64	61.15	38.45	1.59x	60.99	37.89	1.61x
128	98.89	66.26	1.49x	99.21	66.21	1.50x

Model Evaluation

An lm_eval score comparison was not run: the fused path requires the
num_v_heads == 8 * num_k_heads GDN head layout, and no public Qwen3.5
checkpoint ships that config (0.8B is 16/16, 35B-A3B is 16/32, 397B-A17B is
16/64) — on those models the dispatch check never selects the fused kernel, so
default behavior is unchanged. Correctness of the fused kernel against the
Triton path is covered by the numerical-equivalence tests above:
test_fused_gdn_post_conv.py::test_fused_gdn_decode_post_conv_mtp_ratio8
checks kernel outputs and updated recurrent state against
fused_sigmoid_gating_delta_rule_update + rmsnorm_fn across BF16/FP32 state,
ragged acceptance patterns, and TP head counts, and test_gdn_fused_mtp.py
checks the model-path dispatch (pure / mixed / prefill / decode batches)
end-to-end at the layer level.

Essential Elements of an Effective PR Description Checklist

The purpose of the PR, such as "Fix some issue (link existing issues this PR will resolve)".
The test plan, such as providing test command.
The test results, such as pasting the results comparison before and after, or e2e results
(Optional) The necessary documentation update, such as updating supported_models.md and examples for a new model.

claude

Claude Code Review

This pull request is from a fork — automated review is disabled. A repository maintainer can comment @claude review to run a one-time review.

gau-nernst

Added some comments on style and default behavior. The GDN dispatch logic is getting quite complicated, I need more time to review it. Thank you 🙏

gau-nernst · 2026-08-11T01:27:31Z

If this is only used by fused_gdn_decode_kernel.cu, I wonder if it's better to just inline it in the .cu file

gau-nernst · 2026-08-11T01:28:44Z

+    splitting_ops = config.compilation_config.splitting_ops
+    assert splitting_ops is not None
+    assert {
+        "vllm::qwen_gdn_attention_core_fused_norm_packed",
+    } <= set(splitting_ops)


Q: what does this do?

This assertion is a regression guard for the default piecewise-compilation config.

The fused MTP decode path is invoked through a new custom op, vllm::qwen_gdn_attention_core_fused_norm_packed (the packed variant of qwen_gdn_attention_core that also carries the output gate for the fused norm). Like the other GDN/mamba attention-core ops, it mutates the recurrent state cache in place and dispatches on attention metadata at runtime, so it must not be traced into the compiled subgraphs — it has to act as a graph-split boundary for piecewise CUDA graphs. That's why the PR adds it to CompilationConfig._attention_ops (see vllm/config/compilation.py), which is what populates splitting_ops when the engine resolves the default FULL_AND_PIECEWISE cudagraph mode.

test_splitting_ops_dynamic already verifies how splitting_ops gets populated under the default config; the added assertion just pins that the new op is included there by default. Without it, a refactor that dropped the op from _attention_ops would still pass the kernel tests but silently let torch.compile trace through the op, breaking piecewise CUDA graph capture for the fused GDN MTP path.

gau-nernst · 2026-08-11T01:34:16Z

  if(FUSED_KDA_DECODE_ARCHS)
    set(FUSED_KDA_DECODE_SRC
-      "csrc/libtorch_stable/kimi_k3/fused_kda_decode_kernel.cu")
+      "csrc/libtorch_stable/kimi_k3/fused_kda_decode_kernel.cu"
+      "csrc/libtorch_stable/gdn/fused_gdn_decode_kernel.cu")
    set_gencode_flags_for_srcs(
      SRCS "${FUSED_KDA_DECODE_SRC}"
      CUDA_ARCHS "${FUSED_KDA_DECODE_ARCHS}")


nit: since you add the fused GDN decode kernel under this, we probably want to rename it to FUSED_GDN_KDA_DECODE_ARCHS and FUSED_GDN_KDA_DECODE_SRC etc...

Note: currently this also compiles for sm90a and sm120f

Thanks. It now has its own FUSED_GDN_DECODE_ARCHS / FUSED_GDN_DECODE_SRC block and a dedicated VLLM_ENABLE_FUSED_GDN_DECODE define, separate from FUSED_KDA_DECODE_*. The split is necessary rather than cosmetic: the two kernels now build for different arch sets (KDA stays on 9.0a;10.0f;12.0f, while the GDN decode kernel builds for 8.0;8.6;8.9;9.0a;10.0f;12.0f since it only needs cp.async and BF16 math). So each name now matches exactly onw kernel - no shared naming left to rename.

gau-nernst · 2026-08-11T01:35:06Z

+    # Select the Qwen3.5 GDN decode implementation.
+    "VLLM_QWEN3_5_GDN_DECODE_KERNEL": env_with_choices(
+        "VLLM_QWEN3_5_GDN_DECODE_KERNEL",
+        "triton",
+        ["triton", "fused"],
+        case_sensitive=False,
+    ),


Why can't we enable it by default?

I haven't tested it on various datasets yet, it only passes the gsm8k and random data tests.
So, I prefer making it an env option currently.

I'm not familiar with the development guide of vLLM, if you think it's OK to enable it default. I'm happy to do change it.

I think it's fine to enable it by default as long as the constraints are met cc @ZJY0516

I think it's okay if it's always faster

BTW, I prefer the name like "cuda" or "cuda-fused", instead of "fused"

Done, switch to "cuda" instead of "fused" and make it default.

gau-nernst · 2026-08-11T01:38:44Z

+            or vllm_config.model_config.dtype != torch.bfloat16
+            or conv_state_dtype != torch.bfloat16
+            or recurrent_state_dtype not in FUSED_GDN_STATE_DTYPES
+            or not current_platform.is_device_capability_family(100)


At a quick glance, the kernel doesn't seem to use sm100-exclusive features. It should work on sm80 and after? (also further gated by the Cmake configuration below)

thanks, I have extended the constraint to the the compute capability >= 80

gau-nernst · 2026-08-11T01:39:28Z

+            core_attn_out = torch.zeros(
+                (num_tokens, self.num_v_heads // self.tp_size, self.head_v_dim),
+                dtype=hidden_states.dtype,
+                device=hidden_states.device,
+            )


Is torch.zeros() necessary (instead of torch.empty())?

Good question. torch.empty would indeed save some memset overhead, but it's intentionally torch.zeros here, matching the non-fused allocation a few lines below which carries an explicit comment pointing at #28182. That PR tried exactly this swap for the GDN buffers and is still open: with full CUDA graphs the batch is padded, and the GDN core (both the Triton path and this fused kernel, which only writes the rows covered by spec_query_start_loc) fills only the unpadded rows, so torch.empty leaves garbage in the padded tail that then flows through out_proj into the hidden states. The warmup path (attn_metadata is None) doesn't write the buffer at all. I'd prefer to keep the two call sites consitent; if #28182 lands with a safe way to drop the zero-init, both can be switched together.

@ZJY0516 told me the original issue has been fixed (I use torch.empty() for the Kimi-K3-owned KDA layer). But it's fine, just a small detail, we can leave it for future PRs.

gau-nernst · 2026-08-11T01:41:57Z

+    def _output_projection_normalized(
+        self,
+        core_attn_out: torch.Tensor,
+    ) -> torch.Tensor:
+        core_attn_out = core_attn_out.flatten(-2)
+        output, _ = self.out_proj(core_attn_out)
+        return output


nit: just inline this

Done, thanks

ZJY0516 · 2026-08-12T03:41:47Z

/ci run

github-actions · 2026-08-12T03:42:47Z

✅ @Jie-Fang, CI is now available for this PR.

/ci run starts a CI build.
/ci retry retries failed jobs in the CI build for the current PR head. If the current head has no CI build, it starts a new CI build for the current head containing only jobs that failed in the latest earlier CI build for this PR.
/ci cancel cancels scheduled or running CI builds for this PR branch.

github-actions · 2026-08-12T03:42:58Z

✅ Triggered Buildkite CI #83496 for commit 9a812dca34fc.

mergify · 2026-08-12T03:51:09Z

Hi @Jie-Fang, the pre-commit checks have failed. Please run:

uv pip install pre-commit>=4.5.1
pre-commit install
pre-commit run --all-files

Then, commit the changes and push to your branch.

For future commits, pre-commit will run automatically on changed files before each commit.

The fused kernel now handles only pure spec-decode batches; mixed batches fall back to the Triton path. Fusing the mixed case required spec_token_prefix_len metadata and prefix-slicing logic across the core forward for little measured gain. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Jie Fang <jief@nvidia.com>

The kernel only relies on cp.async and BF16 math, both available since SM80. Give it a dedicated arch list and VLLM_ENABLE_FUSED_GDN_DECODE define instead of piggybacking on the fused KDA decode configuration. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Jie Fang <jief@nvidia.com>

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Jie Fang <jief@nvidia.com>

VLLM_QWEN3_5_GDN_DECODE_KERNEL={triton,fused} becomes VLLM_GDN_DECODE_KERNEL={cuda,triton} with cuda as the default: the kernel's constraints are shape- and dtype-based rather than tied to Qwen3.5. When a model or platform doesn't meet them the default falls back to the Triton path; an explicit cuda setting raises with the reason. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Jie Fang <jief@nvidia.com>

The rebase conflict resolution kept upstream's a_spec/b_spec index_select but left the delta-rule call reading the full-batch a/b, undoing the upstream fix for mixed batches. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Jie Fang <jief@nvidia.com>

gau-nernst · 2026-08-13T14:49:58Z

/ci run

github-actions · 2026-08-13T14:50:17Z

✅ Triggered Buildkite CI #83745 for commit 38a1c17991ba.

Conflicts: vllm/envs.py, tests/test_envs.py. Both resolved take-ours. main's legacy `if TYPE_CHECKING:` block and `environment_variables` dict are superseded wholesale by the pydantic BaseSettings tree on this branch; main's conflicting test tail covers `env_with_choices` / `env_list_with_choices`, helpers this branch deleted. Ported from main (b216db3..03a8d0b): - vllm-project#51674 (1be3628), fused CUDA post-conv MTP decode kernel for Qwen3.5 GDN. VLLM_GDN_DECODE_KERNEL -> `gdn_decode_kernel: Literal["cuda", "triton"] = "cuda"` in UsageSettings, next to `enable_fla_packed_recurrent_decode`. main passes `case_sensitive=False`, so a bare Literal would reject the "CUDA" main accepts; added `_lower_gdn_decode_kernel`, mirroring `_lower_mm_hasher`. Deliberately no strip: main's env_with_choices does not strip, so " cuda " must keep raising. Caller qwen_gdn_linear_attn.py:492. - vllm-project#50487 (03a8d0b), Kimi K3 DFlash aux state. VLLM_KIMI_K3_AUX_ATTN_RES_STREAM -> `kimi_k3_aux_attn_res_stream: bool = False` in QuantSettings, between `kimi_k3_shard_sp_shared_expert` and `kimi_k3_gemm_rs` to match main's ordering. main's `bool(int(getenv(...)))` needs no validator. Caller kimi_k3/nvidia/model.py:1218. Both vars have already-merged callers, so both ports are mandatory. Neither was added to main's `ignored_factors`, so both remain compile factors on the branch and carry no `compile_factor: False` marker; verified at runtime. Not ported: main's `test_gdn_decode_kernel_env`, which exercises the deleted `env_with_choices` helper through the back-compat shim. No branch-flavored replacement was added (user decision); the lowercase coercion was verified by hand instead. Nothing else dropped, both commits accounted for. Parity check: 294 branch fields vs 295 main runtime entries, sole difference VLLM_TRITON_ATTN_USE_TD, the known shim divergence, re-confirmed byte-identical in base and theirs. Tests: 54 passed across tests/test_envs.py, tests/test_envs_pydantic.py and tests/docs/test_env_vars_gen.py; `pre-commit run --files vllm/envs.py tests/test_envs.py` clean. The GPU-only consumer suites (tests/kernels/mamba/test_gdn_fused_mtp.py, tests/models/kimi_k3/test_eagle3.py) were not run here. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Signed-off-by: Vinay Damodaran <vrdn@hey.com>

…_conv_mtp declared under KDA ifdef The declaration sat inside VLLM_ENABLE_FUSED_KDA_DECODE while its binding (torch_bindings.cpp:811) is guarded by VLLM_ENABLE_FUSED_GDN_DECODE — any build with GDN archs but no KDA archs (sm_89 wheel builder) fails 'not declared in this scope'. Declaration moved to its own GDN guard.

Upstream made it prefill-lazy; the ReplaySSM cursor-reset and ring- coherence blocks read it on pure-(spec-)decode batches (incl. cudagraph capture), raising UnboundLocalError at capture time. Restore the pre-vllm-project#51674 unconditional compute (one tensor sub). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

… GDN (vllm-project#51674) Signed-off-by: Jie Fang <jief@nvidia.com> Co-authored-by: OpenAI Codex <codex@openai.com> Co-authored-by: Claude Fable 5 <noreply@anthropic.com>

… GDN (vllm-project#51674) Signed-off-by: Jie Fang <jief@nvidia.com> Co-authored-by: OpenAI Codex <codex@openai.com> Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Wyett <wyettzeng@gmail.com>

… GDN (vllm-project#51674) Signed-off-by: Jie Fang <jief@nvidia.com> Co-authored-by: OpenAI Codex <codex@openai.com> Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Zhu, Zufang <zufang.zhu@intel.com>

The fused GDN decode kernel (vllm-project#51674) only served the MTP speculative path; plain decode still ran the Triton packed recurrent kernel plus a separate gated-RMSNorm launch per layer. Invoke the fused kernel with draft width 1 for pure-decode batches: each request loads its recurrent state slot, advances it in place, and gets the gated-RMSNorm epilogue in the same launch. Per-layer eager cost drops ~10x (127us -> 12us, L40S, B<=8); Qwen3.8-27B TP2 enforce-eager serving improves 15.3% output throughput / -13.6% TPOT. Neutral under FULL CUDA graphs where launch overhead is already amortized (graph-replay GPU time is on par with the Triton path). Signed-off-by: Cheng Rui <286040359@qq.com> Co-authored-by: Claude <noreply@anthropic.com>

… GDN (vllm-project#51674) Signed-off-by: Jie Fang <jief@nvidia.com> Co-authored-by: OpenAI Codex <codex@openai.com> Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: khushali9 <khushali.desai9@gmail.com>

The fused GDN decode kernel (vllm-project#51674) only served the MTP speculative path; plain decode still ran the Triton packed recurrent kernel plus a separate gated-RMSNorm launch per layer. Invoke the fused kernel with draft width 1 for pure-decode batches: each request loads its recurrent state slot, advances it in place, and gets the gated-RMSNorm epilogue in the same launch. Per-layer eager cost drops ~10x (127us -> 12us, L40S, B<=8); Qwen3.8-27B TP2 enforce-eager serving improves 15.3% output throughput / -13.6% TPOT. Neutral under FULL CUDA graphs where launch overhead is already amortized (graph-replay GPU time is on par with the Triton path). Signed-off-by: Cheng Rui <286040359@qq.com> Co-authored-by: Claude <noreply@anthropic.com>

Jie-Fang requested review from AndreasKaratzas, Harry-Chen, LucasWilkinson, ProExpertProg, WoosukKwon, ZJY0516, houseroad, mgoin, robertgshaw2-redhat, tdoublep, tlrmchlsmth, vadiklyutiy, yewentao256, youkaichao and zyongye as code owners August 10, 2026 12:49

claude Bot reviewed Aug 10, 2026

View reviewed changes

mergify Bot added ci/build qwen Related to Qwen models nvidia labels Aug 10, 2026

github-project-automation Bot added this to NVIDIA Aug 10, 2026

gau-nernst reviewed Aug 11, 2026

View reviewed changes

shizhehe mentioned this pull request Aug 11, 2026

[Bug]: Qwen3.5 GDN CUDA kernels promote BF16 Q/K normalization to FP32 #51779

Open

ZJY0516 mentioned this pull request Aug 11, 2026

[Bugfix] Align Qwen GDN gates with speculative tokens #51812

Merged

gau-nernst approved these changes Aug 12, 2026

View reviewed changes

ZJY0516 added the ready ONLY add when PR is ready to merge/full CI is needed label Aug 12, 2026

feednetinfra mentioned this pull request Aug 12, 2026

[Perf][Qwen3.5] Avoid GDN decode gate copies #51954

Open

Jie-Fang and others added 6 commits August 13, 2026 07:37

Inline the fused GDN output projection helper

dc85d7c

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Jie Fang <jief@nvidia.com>

Apply clang-format to the fused GDN decode kernel

7a99fc7

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: Jie Fang <jief@nvidia.com>

Jie-Fang force-pushed the gdn-decode-mtp branch from bf57415 to 38a1c17 Compare August 13, 2026 14:44

ZJY0516 approved these changes Aug 14, 2026

View reviewed changes

github-project-automation Bot moved this to Ready in NVIDIA Aug 14, 2026

ZJY0516 merged commit 1be3628 into vllm-project:main Aug 14, 2026
250 checks passed

github-project-automation Bot moved this from Ready to Done in NVIDIA Aug 14, 2026

huijunwu mentioned this pull request Aug 16, 2026

Fix: move fused_gdn_decode_post_conv_mtp declaration to GDN guard huijunwu/vllm#1

Draft

BabyDrangoner mentioned this pull request Aug 16, 2026

[Kernel][Perf] Support Qwen head ratios in fused GDN MTP #52539

Merged

prakharPant mentioned this pull request Aug 22, 2026

fix(build): correct preprocessor guard for GDN decode to fix Ampere c… #52743

Merged

4 tasks

positive666 mentioned this pull request Aug 23, 2026

[Kernel][Perf] Route non-spec GDN decode through the fused CUDA kernel #53463

Open

mayuyuace mentioned this pull request Aug 27, 2026

[Fix] Fix sm8x build failure: move fused_gdn_decode_post_conv_mtp to correct ifdef guard #54008

Closed

chiptoe-svg mentioned this pull request Sep 5, 2026

[Bug] FlashInfer + MTP speculative decoding crashes on SM121 (DGX Spark) with GQA=16 model #37754

Open

jahnclawdmonet mentioned this pull request Sep 8, 2026

[Bug]: MTP + FlashInfer long-context CUDA IMA / Xid 31 still on v0.27.1 (Qwen3.8-27B-NVFP4, RTX 5090, max-num-seqs=1) #55775

Open

1 task

Uh oh!

Uh oh!

Conversation

Jie-Fang commented Aug 10, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Purpose

Test Plan

Test Result

Microbenchmark

Model Evaluation

Uh oh!

claude Bot left a comment

Choose a reason for hiding this comment

Claude Code Review

Uh oh!

gau-nernst left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Jie-Fang Aug 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

ZJY0516 commented Aug 12, 2026

Uh oh!

github-actions Bot commented Aug 12, 2026

Uh oh!

github-actions Bot commented Aug 12, 2026

Uh oh!

mergify Bot commented Aug 12, 2026

Uh oh!

gau-nernst commented Aug 13, 2026

Uh oh!

github-actions Bot commented Aug 13, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Jie-Fang commented Aug 10, 2026 •

edited

Loading

Jie-Fang Aug 11, 2026 •

edited

Loading