These 15 commits optimize two hot subsystems for Qwen3.5-35B-A3B — GDN (the gated-delta-net linear attention layer) and MoE — plus KV-cache handling and build fixes.

GDN linear attention

373565d tilelang gated RMSNorm
The GDN output gating (RMSNorm + SiLU gate) was a chain of elementwise kernels; replaced with a single fused tilelang kernel.

7fa8fe4 fused zba split
The GDN forward splits z/b/a out of the packed projection. Replaces split + reshape + two contiguous copies with one kernel that slices contiguous z/b/a directly.

9e61f28 mate decode
Single-token GDN decode now runs mate's fp32 VK recurrent kernel (was the slower Triton fused-recurrent update), with automatic fallback.

f046133 tilelang prefill causal conv1d
The causal 1D convolution in GDN (prefill) moved to a tilelang kernel.

98a253d strided q/k/v + fused gating
Feeds mate strided q/k/v views instead of materializing contiguous copies (zero-copy), and fuses the log-space GDN gating (g, beta) into one kernel — replacing ~9 elementwise kernels (ported from SGLang).

b7d72d4 drop D2H + 3D→2D
Removes a per-call device→host sync in conv1d (reads a CPU-side value already carried on the metadata); routes the q/k GemmaRMSNorm through the kernel for 3D inputs (which previously fell back to the slow path).

fd9ad37 mate output buffer + int32
Lets mate write straight into the caller's buffer, saving one D2D copy; switches cu_seqlens from int64 to int32 (mate specializes on int32 internally).

MoE

b03acb0 generalize DeepGEMM MoE prefill
Extends the DeepGEMM fused-prefill path from DeepSeek-V4-only to any qualifying block-FP8 MoE ( scales, etc.). Default-on, gated by VLLM_MUSA_MOE_DEEPGEMM_PREFILL. This is the main TTFTlever.

8af8300 shared-expert fold
Qwen3.5's shared expert ran as a separate serial MLP, adding a chain of kernels per layer. Folds it into the routed grouped GEMM as an always-selected extra expert (slot 256), producing the result in one kernel. Handles both unquantized and FP8 weights; disabled under expert parallelism.

1f369d9 row-tiled fused silu+mul+FP8 quant
Fuses the MoE silu_and_mul and per-token-group FP8 quantization into one row-tiled kernel.

5f8dcf6 per-token-group vec quant
group-128 bf16/fp16 activation quant now uses a register-resident vec kernel (was routing through the slower upstream path); caches a dummy tensor to avoid a per-call torch.empty.

KV cache

6558174 + a4b544c separate mamba state pool
Gives the mamba (GDN) state its own contiguous BlockPool with a private block-id space, so mate's GDN decode updates the recurrent state in place (no gather/scatter). This also fixes an overlap bug where attention KV and mamba state shared one memory tensor (garbage decode output under high concurrency), so it is defaulted on.

QK-norm + positional encoding

7732461 fused QK-RMSNorm + MRoPE
For interleaved-MRoPE gated attention, fuses QK-RMSNorm and MRoPE into one kernel. These modelsry path and emit a long tail of elementwise kernels. Dispatched ahead of the IR op, which staysas the fallback. Gated by VLLM_MUSA_FUSED_QK_MROPE.

Build fixes

a90b1fc rebase two build patches
0060 and 0085 no longer matched their surrounding series after upstream changes, breaking a clean install; their context was re-aligned.

Perf test

Qwen3.5-35B-A3B-FP8 · TP1 · FULL_DECODE_ONLY · 8 prompts (~2.5k in / 1.5k out)

============ Serving Benchmark Result ============
Successful requests:                     8
Failed requests:                         0
Maximum request concurrency:             64
Benchmark duration (s):                  57.55
Total input tokens:                      20083
Total generated tokens:                  12000
Request throughput (req/s):              0.14
Output token throughput (tok/s):         208.53
Peak output token throughput (tok/s):    216.00
Peak concurrent requests:                8.00
Total token throughput (tok/s):          557.51
Mean TTFT (ms):                          913.54
Median TTFT (ms):                        957.80
P99 TTFT (ms):                           1133.53
Mean TPOT (ms):                          37.76
Median TPOT (ms):                        37.73
P99 TPOT (ms):                           38.03
Mean ITL (ms):                           37.73
Median ITL (ms):                         37.45
P99 ITL (ms):                            41.27
==================================================

============ Serving Benchmark Result ============
Successful requests:                     8
Failed requests:                         0
Maximum request concurrency:             64
Benchmark duration (s):                  30.29
Total input tokens:                      20083
Total generated tokens:                  12000
Request throughput (req/s):              0.26
Output token throughput (tok/s):         396.12
Peak output token throughput (tok/s):    448.00
Peak concurrent requests:                8.00
Total token throughput (tok/s):          1059.05
Mean TTFT (ms):                          536.78
Median TTFT (ms):                        557.22
P99 TTFT (ms):                           680.83
Mean TPOT (ms):                          19.84
Median TPOT (ms):                        19.82
P99 TPOT (ms):                           19.99
Mean ITL (ms):                           19.82
Median ITL (ms):                         19.92
P99 ITL (ms):                            20.72
==================================================