Release Highlights

This release completes NVIDIA Rubin (SM107) support, brings whole-layer expert-parallel MoE to Hopper, adds decode paths for DeepSeek-V4 sparse attention and Kimi K3 linear attention, and broadens low-precision MoE coverage with MXFP4 on Blackwell RTX PRO and DGX Spark and weight-only W4A16 on B200 and B300.

SM107 (Rubin) support

FlashInfer 0.6.18 completes support for NVIDIA Rubin (SM107), begun in 0.6.16. Rubin devices now dispatch through the same unified APIs as Blackwell — attention, GEMM, MoE, and quantization.

Attention gains trtllm-gen FMHA for SM107, including sparse compression and FP16 softmax. The PrimTS attention path accepts Rubin as well. On the GEMM side, batched and low-latency GEMM both run on Rubin, and CUTLASS NVFP4 SVDQuant is enabled. A CuTe-DSL kernel family specialized for Rubin ships alongside them. trtllm-gen MoE now selects valid tactics on Rubin, and router GEMM and topk_varlen accept SM107.

• #4526 feat: CuTe DSL kernels for Rubin (SM107) and batched FP8 GEMM for Blackwell
• #4596 feat: trtllm-gen FMHA features for sm107 (spcompress, fp16softmax)
• #4509 gemm: enable CUTLASS NVFP4 SVDQuant on SM107
• #4710 fix: support SM107 (Rubin) in topk_varlen, FlashKDA, and router GEMM capability gates
• #4792 [release-v0.6.18] Five SM107 (Rubin) fixes: 133 CI failures

Expert-parallel MoE runs whole-layer on Hopper

The moe_ep mega-kernel stack was Blackwell-only. Two Hopper FP8 backends, Sm90PullFp8MegaMoeConfig and Sm90PushFp8MegaMoeConfig, now let Hopper deployments run dispatch, FC1, SwiGLU, FC2, and combine as one fused layer behind the existing MoEEpLayer, instead of composing an NCCL all-to-all with a local fused-MoE operator. The pull backend reaches 562 TFLOPS/rank at a 384-expert DeepSeek-class geometry. The push backend supports CUDA Graph capture, and its opt-in fused FC1 epilogue drops an approximately 1 GiB per-rank activation buffer at the DeepSeek-V3 EP8 shape. Its grouped GEMM requires CUDA Toolkit 12.8 or newer.

• #4113 feat(moe_ep): SM90 (Hopper) pull-style FP8 mega-MoE backend
• #4069 feat(moe_ep): add SM90 push FP8 mega-MoE backend for Hopper
• #4449 refactor(moe_ep): taxonomy/provenance restructure; incorporate SM90 push-style FP8 backend; sync CuTe-DSL 4.7 quant-staging fix

HCA decode backend and top-K selection for DeepSeek-V4

Sparse attention picks the top-K KV positions per request on every decode step, then attends over the compressed cache; both halves now have dedicated paths. FP8 Heavily Compressed Attention (HCA) arrives for SM100/SM103 via trtllm_batch_decode_sparse_mla_dsv4(..., backend="cute-dsl"), taking arbitrary sliding-window row order including ring rotation and wraparound while keeping the compressed cache paged. The new flashinfer.top_k_varlen handles ragged batches through a Blackwell radix kernel, a guess-verify-refine kernel that warm-starts from the previous step’s indices, and a CUTLASS fallback for any GPU. SM120/121 also picks up top-k 192 and 256.

• #3943 Add CuTe DSL HCA backend for DeepSeek V4 sparse MLA
• #3901 feat(topk): Add top_k_varlen with GVR and radix backends for sparse-attention KV selection
• #4380 feat(sm120): consolidate DSV4 sparse MLA top-k 192/256 support

Kimi K3 decode fuses into one Blackwell kernel

flashinfer.fused_kda_decode folds Kimi K3’s width-four depthwise causal convolution, SiLU, recurrent Kimi Delta Attention update, and gated RMSNorm into a single SM100 launch, covering the production head_dim 128 shapes at 12, 24, 48, and 96 heads and updating the convolution cache and FP32 state in place. On B200 under CUDA Graphs it is 1.33x the vLLM fused kernel at one row (table geomean 1.13x). A T=1 fast path inside recurrent_kda and shared SM100-family recurrent kernels round out the KDA stack.

• #4243 feat(kda): add fused Kimi K3 decode kernel
• #4417 feat(kda): add packed-input CuTe decode kernel
• #4562 feat(cake_kda): add optimized H12 packed decode across SM100 family

MXFP4 MoE and video sparse attention on Blackwell RTX PRO

MXFP4 checkpoints run natively on SM120/121: b12x_fused_moe and B12xMoEWrapper accept quant_mode="mxfp4" across the existing fused schedules including CUDA Graph reuse, the b12x dense GEMM gains the matching path, and on an RTX PRO 6000 Blackwell Server Edition MXFP4 tracks NVFP4 against a strict quantized reference. Video Sparse Attention, until now datacenter-Blackwell only, reaches these parts through a vsa_sm120_blk64 backend on BlockSparseAttentionWrapper. Gemma 4 gains asymmetric VO-split NVFP4 paged prefill on SM120/121.

• #4290 feat(sm120): add native MXFP4 W4A4 fused MoE
• #4479 feat(b12x): support MXFP4 dense GEMM on SM120
• #4259 integrate SM120 VSA (Video Sparse Attention) block-sparse backend
• #3684 feat(attention): asymmetric VO-split NVFP4 paged prefill (qk=512/vo=256) for Gemma-4 on SM120/121

W4A16 MoE and dense GEMM extend to B200 and B300

Weight-only NVFP4 against BF16 activations, which 0.6.14 shipped for SM12x, now runs on the SM100 family: CuteDslMoEWrapper and cute_dsl_fused_moe_nvfp4 accept quant_mode="w4a16", decoding weights to BF16 inside the kernel so no separate activation-quantization or repack launch is needed, and mm_bf16_fp4 gains a dedicated SM100/SM103 kernel. Consuming BF16 directly pays off where MoE decode is memory-bound: at a DeepSeek EP8 shape on B200, W4A16 is 1.50x the W4A4 baseline at one token; W4A4 still wins large-batch prefill. GeGLU-tanh and SiTU are supported.

• #4048 feat(moe): support BF16 activations in CuTe DSL NVFP4 MoE
• #4466 feat: sm100 cute_dsl w4a16 gemm
• #4394 feat(moe): support GeGLU and SiTU in W4A16

Unified MoE API adds shared experts, MXINT4, and CUTLASS runners

Shared experts now work through the unified API rather than low-level kernel entry points, via ExpertConfig.num_fused_shared_experts on the block-FP8 and FP4 runners, completing in the unified API what 0.6.15 and 0.6.17 added for FP8 and FP4. The API also gains MXINT4, CUTLASS BF16 and W4A16 runners on SM90, packed per-tensor FP8 routing, BF16 FromLogits routing, and TopKSigmoid.

• #4326 Fused shared experts support via the unified API
• #4320 feat(moe): enable MxInt4 in unified API
• #4328 feat(moe): add unified CUTLASS BF16 and W4A16 runners
• #4227 feat(moe): add packed per-tensor FP8 and BF16 FromLogits routing to unified MoE API
• #4404 feat(moe): add TopKSigmoid routing method (TopK -> Sigmoid)

Fused MNNVL all-reduce tail for tensor-parallel MoE

allreduce_fusion gains a BF16 Blackwell CuTe-DSL backend for the latency-critical tail of tensor-parallel MoE layers, fusing all-reduce, residual add, and RMSNorm, optionally preceded by MoE finalize and shared-expert add, over MNNVL/NVLink multicast. One backend spans decode to prefill by switching protocol with token count, with initial profiles targeting GB300 on TP8 and TP16.

• #4358 feat(comm): add Blackwell MNNVL CuTe DSL all-reduce fusion backend

Smaller JIT-cache wheels; SM75 and single-request FA2 are JIT-only

The flashinfer-jit-cache wheels no longer ship precompiled kernels for SM75 (Turing). Those GPUs still run; the kernels compile on first use. CUDA 13 AArch64 wheels also drop native SM121a cubins (DGX Spark keeps running via SM120 family cubins), and the single-request single_decode_with_kv_cache / single_prefill_with_kv_cache FA2 modules are no longer AOT-prebuilt — those APIs still JIT. Fatbins use size-oriented compression.

• #4757 build(jit): reduce JIT-cache wheel size
• #4760 chore(aot): exclude single prefill/decode modules from jit-cache prebuilds

This release brings MoE expert parallelism into serving engines, refreshes the Blackwell SM12x fused-MoE kernels with an FP4 accuracy fix, extends the unified MoE API to MXFP4 W4A8/W4A16 and shared experts, and adds decode coverage for Kimi K3 MLA and MiniMax-M3 sparse attention on vLLM.

MoE expert parallelism production-ready in vLLM

The MegaMoE path in flashinfer.moe_ep is ready for serving engines: full CUDA-graph capture and replay, a fused single-launch quantize-and-stage hot path, prequantized weight packs, symmetric-buffer workspaces pooled across layers, and a persistent knob cache that resolves tuned knobs by lookup, so production sessions run with no in-engine autotuning. An opt-in fault-tolerance rank mask, available over both NCCL-EP and NIXL-EP, masks and skips a peer that times out during dispatch or combine, keeping the job alive when a rank dies. Deployment gets simpler too: the CuTe-DSL runtime floor returns to 4.5.2, the version vLLM 0.25.1 pins, and a new BootstrapConfig.device lets the host framework pin each worker’s CUDA device.

• #4079 feat(moe_ep): MegaMoE framework integration ready: CUDA graph support, fused quant+stage launch, persistent knob cache, and prequantized weight packs
• #4183 feat(moe_ep): fault-tolerance rank mask (NCCL-EP + NIXL-EP)
• #4101 perf(moe_ep): CuTe-DSL 4.5.2 mainloop WAR — drop the 4.6.1 runtime floor
• #4348 fix(moe_ep): let the host framework pin the CUDA device via Bootstrap

Blackwell SM12x fused MoE refreshed, with an FP4 accuracy fix

W4A4 serving on DGX Spark and RTX PRO parts (GB10, SM120/SM121) now delivers the output quality its benchmark scores imply, with two NVFP4 quantization bugs fixed and a new input_global_scale that lets integrators pass a checkpoint’s weight scale directly. The SM12x fused-MoE families are synced to current b12x: the NVFP4 W4A4 backend reaches kernel parity across decode and prefill, and the W4A16 family adds cooperative persistent launches, a tensor-core decode path for small batches, and shape-stable route packing that keeps decode batch-size changes recompile-free.

• #3932 fix(b12x): correct fp4 quantization numerics and add input_global_scale to decouple weight and activation scales
• #4285 perf(moe): sync SM12x NVFP4 fused-MoE kernels to b12x HEAD
• #4255 feat(moe): sync SM12x W4A16 fused MoE family to b12x HEAD
• #4253 feat(gemm): sync mm_fp4 SM120 NVFP4 dense GEMM kernel to b12x HEAD
• #4130 feat(sm120): fused MoE (SwiGLU) via moe_gemm is_gated for cute SM120 groupwise GEMM

Unified MoE API adds MXFP4 W4A8 and W4A16, shared experts, and SiTU

TRTLLM-gen MXFP4 weights now run through the unified MoELayer API against both MXFP8 activations (W4A8) and BF16 activations (W4A16), completing the FP8 series begun in 0.6.16, alongside per-tensor routed FP8. Routing adds an unpacked pre-routed FP4 mode that accepts topk_ids and topk_weights as separate contiguous tensors. TRTLLM-gen FP4 MoE also gains shared-expert fusion, extending to FP4 what 0.6.15 added for FP8, plus SiTU activation for MXFP4 x MXFP8 and NVFP4 x NVFP4.

• #4159 feat(moe): add TRTLLM MXFP4 W4A8 and W4A16 unified API support
• #4088 feat: close feature gap by wiring up per-tensor routed FP8 fused-moe
• #4104 feat(moe): add unified unpacked pre-routed FP4 mode
• #4239 feat(moe): support shared expert fusion for trtllm-gen fp4 moe
• #4180 [feat] Add SITU trtllmgen MOE

Kimi K3 MLA decode, and MiniMax-M3 sparse attention under vLLM

Blackwell decode now covers Kimi K3’s MLA geometry — 96 global query heads against one KV head, TP-local head counts down to 6, speculative query lengths up to 8, and context parallelism for long contexts — by packing query-token and query-head rows into shared CuTe-DSL tiles, adding compact variable-length Q, and extending TRTLLM-gen dense and sparse MLA to non-power-of-two head counts. Sparse MLA decode also serves the no-rotary-tail shape (kv_lora_rank=512, qk_rope_head_dim=0) natively. MiniMax Sparse Attention accepts vLLM’s packed paged KV layout for MiniMax-M3, unblocking the vLLM integration on SM120/SM121, with lower per-call overhead on paged decode.

• #4178 feat(mla): support packed low-head and variable-Q decode
• #4108 feat: support native qk_rope_head_dim=0 sparse MLA decode in trtllm-gen
• #4039 feat(msa): accept K/V views split from a packed paged KV cache
• #4324 MSA decode path improvements

Ulysses sequence parallelism for long-context and video diffusion

Ulysses sequence parallelism gets its head-scatter / sequence-gather all-to-all as a public API, for video diffusion transformers and other long-sequence attention workloads. A fused NVLink-P2P kernel folds the layout permutation directly into the cross-GPU writes over CUDA IPC for a single coalesced push, with automatic NCCL fallback when P2P is unavailable.

• #3820 Feat/ulysses p2p a2a
• #4240 docs: improve Ulysses communicator and MoE EP docs

This release delivers the MegaMoE kernels promised for expert parallelism in 0.6.15 and extends Blackwell model coverage — MiniMax-M3 sparse attention on Blackwell RTX and DGX Spark, and a unified MoE API that now executes per-tensor FP8, block-scale FP8, and B12x NVFP4/W4A16 — plus distributed serving under Confidential Computing with multicast-free all-reduce fusion. Under the hood, an on-disk JIT cache for CuTe-DSL kernels cuts cold-start compilation and shrinks the install.

⚠️ FlashInfer 0.6.16.post2 has picked up the latest tvm-ffi compatibility fix.

⚠️ FlashInfer 0.6.16.post4 adds a Python 3.10 compatibility hotfix.

We recommend upgrading to the latest 0.6.16 post-release.

MegaMoE kernels land in expert parallelism (moe_ep)

The MoEEpLayer now unifies the split and mega execution paths under one entry point and adds three mega backends — deep_gemm_mega (FP8/FP4), nvfp4_cutedsl, and mxfp8_cutedsl — each fusing expert-parallel communication with the local MoE in a single symmetric-memory kernel (Blackwell SM100+, NVSHMEM). In a microbenchmark, on a GB200 node (EP=4, DeepSeek-V3-like geometry) the tuned CuTeDSL NVFP4 backend with in-register FP4 combine reaches up to 1.89× the throughput of deep_gemm_mega at 8192 tokens/rank. NIXL-EP transport gaps and combine deadlocks blocking the vLLM Fleet/Handle adapters are also closed.

• #3852 feat: MegaMoE Kernel integration in moe_ep
• #3980 feat(moe_ep): restructure MegaMoE kernel sources into kernel_src + latest CuTeDSL MegaMoE kernels with improved performance
• #4075 feat(moe_ep): close the nixl_ep transport gaps blocking the vLLM Fleet/Handle adapters
• #4139 fix(moe_ep): fix two NIXL-EP combine deadlocks under real serving load

MiniMax Sparse Attention (MSA) on Blackwell RTX and DGX Spark

MiniMax-M3’s MSA — a proxy/top-k indexer plus sparse prefill, decode, and combine — now runs on Blackwell RTX and DGX Spark (SM120/121) GPUs. The tensor-core kernels are rebuilt on SM12x warp-level mma.sync, top-k and combine are rewritten in CuTe-DSL, and an optional NVFP4 indexer is available, with prefill/decode accepting FP8 or NVFP4 KV (paged or flat). New APIs live under flashinfer.msa_ops.

• #3655 feat(attn): Enable MiniMax Sparse Attention (MSA) for Consumer Blackwell GPUs (SM120/121)

XQA decode adds sliding-window, attention sinks, and ragged Q for speculative decode

The XQA decode kernel — used on Blackwell RTX and DGX Spark (SM120/121) for models with attention sinks — expands coverage for sliding-window attention and ragged Q, across causal and non-causal draft-block mask modes and combinations of them. Ragged Q lets each request in a batch verify a different number of draft tokens, and sliding-window masking is now computed per draft-token row, making the kernel useful for speculative-decoding workloads across models on SM120/121.

• #4137 feat(xqa): ragged Q and per-row sliding-window masking for speculative decode

Unified MoE API reaches parity with the legacy FP8 and NVFP4 paths

The unified MoELayer API continues to expand quantization formats to reach parity with legacy flat APIs: TRTLLM per-tensor FP8 (SM100/SM103, with Llama4 routing-scale-on-input), DeepSeek FP8 and MXFP8 block-scale (SM100/SM103), and SM120/SM121 B12x NVFP4 and W4A16 backends. In-kernel routing (FromLogits) is now wired through the unified API and fuzzer, so precomputed and in-kernel routing share one path with CUDA-graph and autotuning coverage.

• #4091 feat(moe): add unified per-tensor FP8 support
• #4026 feat(moe): add unified block-scale FP8 support
• #3983 feat(moe): add B12x backends to unified MoE API and add corresponding tests
• #3892 feat(moe): wire in-kernel routing (FromLogits) into unified MoE API + routing axes in fuzzer

Confidential Computing: multicast-free all-reduce fusion and FP8 AllReduce

The TRT-LLM AllReduce-fusion workspace now allocates a multicast-free IPC workspace when NVIDIA Confidential Computing is detected (is_confidential_compute(), overridable via FLASHINFER_CONFIDENTIAL_COMPUTE), so one-shot Lamport and two-shot sync fusion run under CC where cuMulticast setup otherwise fails. Separately, a new flashinfer.comm.quantized_all_reduce() halves AllReduce transfer volume by quantizing activations to FP8 before P2P transfer over symmetric memory (SM90+, NVSwitch).

• #3993 comm: multicast-free (IPC) workspace for TRT-LLM AR fusion under Confidential Computing
• #3759 feat(comm): Add FP8 quantized two-shot AllReduce via symmetric memory

Faster JIT cold-start and smaller install

CuTe-DSL kernels now persist to an on-disk cache (JitSpec gains a JitSpecCuteDsl backend) and reload via JITLink in about 3–30 ms instead of recompiling in every new process; mm_fp4 autotuning additionally compiles tactics in parallel and reuses the shared disk cache, cutting autotune wall time. Pruning architecture gencode that dispatch can never load removes roughly 1.6 GB of installed size from the CUDA-13 aarch64 JIT-cache wheel.

• #3874 feat(jit): JitSpec ABC + disk cache for JIT-compiled CuTe-DSL kernels
• #4029 perf(gemm): Improve mm_fp4 cute-dsl autotune time via disk-cache and parallel compilation
• #3947 perf(jit): drop dead SM12x gencode from SM10x-serving modules
• #4073 perf(jit): prune arch gencode that dispatch can never load

This release ships Expert Parallelism (moe_ep) in the default install and extends the TRTLLM-GEN MoE stack for large-model serving. It broadens Blackwell model coverage — Gemma 4 / MiniMax-M3 MoE on consumer / DGX-Spark SM12x, DeepSeek-class MLA decode on B300, and context-parallel GDN on SM120 — and adds Video Sparse Attention plus CUDA-graph-safe FP8 all-reduce fusion for distributed inference.

Unified MoE API with Expert Parallelism (Experimental API to try!)

FlashInfer’s unified MoE compute API is now wired into expert parallelism. A new flashinfer.moe_ep.MoEEpLayer runs one MoE layer split across ranks as dispatch → per-expert grouped GEMM → combine, over pluggable transport (NCCL-EP via nccl.ep/nccl4py, and NIXL-EP). The expert GEMM reuses the unified flashinfer.fused_moe.MoELayer as a pure per-expert grouped GEMM — routing lives in dispatch/combine. moe_ep is part of the default install (CUDA 13+), with vLLM-facing APIs and checkpoint-safe MoE all-to-all graph VAs. This EP API is also the integration surface for MegaMoE kernels in upcoming releases — try it out and let us know how it works for you.

• #3686 MoE-EP: wire unified MoE compute into NCCL-EP / NIXL-EP expert parallel (LL + HT)
• #3821 feat: make moe_ep (EP) part of the default install; drop nccl submodule
• #3813 Feat/vllm moe ep api
• #3727 feat(mnnvl): preserve MoE all-to-all graph VAs across checkpoint restore
• #3591 test(moe): align EP-offset tests and docs with global expert ids (gh #3547)

Gemma 4 and MiniMax-M3 NVFP4 MoE now run on Blackwell SM12x

Gemma 4 and MiniMax-M3 NVFP4 MoE now run on Blackwell SM12x (consumer / DGX Spark), enabled by two new NVFP4 MoE activation functions — gelu_tanh and swiglu_oai — added to the SM12x MoE path.

• #3744 feat(moe): add gelu_tanh and swiglu_oai activations to b12x NVFP4 MoE for SM12x
• #3737 Feat/cutedsl moe swiglu oai and use activation type to pass in

TRTLLM-GEN MoE adds shared experts, DeepSeek-V4 routing, and low-latency FP8

TRTLLM-GEN MoE gains shared-expert fusion for FP8 paths, hash-based DeepSeek-V4 routing (hash_topk), and a fused FP8 blockwise megakernel that cuts latency for small batches (BS ≤ 8).

• #2625 feat: Fuse shared experts into trtllm_gen moe (fp8)
• #3561 feat(moe): add DSv4 hash-based MoE routing (MOE-01-HASH)
• #3424 Fused FP8 blockwise MoE megakernel For BS <= 8

Video Sparse Attention and DeepSeek MLA decode on Blackwell B300

Video Sparse Attention (VSA) is now integrated into the block-sparse attention API, bringing efficient long-context attention for video diffusion models to FlashInfer. DeepSeek-class MLA decode extends onto Blackwell B300 (SM103) with cluster-aware CUTLASS split_kv, and CuTe-DSL GQA decode adds sliding-window and attention-sink masking for newer attention variants.

• #3250 integrate VSA (Video Sparse Attention) into block-sparse API
• #3888 fix(mla): enable cute-dsl/auto on SM103 + cluster-aware cutlass split_kv
• #3794 feat(cute-dsl): sliding window and attention sinks for GQA decode (non-paged + paged)

Linear attention: context-parallel GDN on SM120 and faster KDA decode

Gated delta-rule (GDN) linear attention adds SM120 context-parallel delta rule support, extending the CuTe-DSL GDN rewrite from 0.6.14 onto Blackwell SM120 prefill, alongside SM90 context-parallel prefill optimizations. KDA recurrent-decode kernels are also optimized for lower decode latency.

• #3659 feat: add sm120 cp delta rule
• #3788 perf: optimize sm90 cp delta rule
• #3766 perf(kda): optimize recurrent decode kernels

Distributed comm: FP8 fusion and CUDA-graph checkpoint restore

Dynamic per-token FP8 quantization fuses into allreduce + residual + RMSNorm (TRT-LLM and MNNVL backends, CUDA-graph safe). All-reduce workspaces support checkpoint_prepare / checkpoint_restore so physical backing can be released and remapped at stable VAs across serving checkpoints.

• #3421 feat: add dynamic FP8 allreduce RMSNorm fusion
• #3745 feat(comm): preserve all-reduce graph VAs across checkpoint restore

This release pushes FlashInfer’s coverage onto Blackwell RTX PRO and DGX Spark silicon, lands the CuTe-DSL rewrite of the gated-delta-rule (GDN) kernels behind Qwen3.5/3.6, and makes those kernels production-ready inside vLLM.

⚠️ Manual intervention needed: install instructions changed

Please modify your install instructions as you bump the version to 0.6.14.

pip install flashinfer-python

pip install flashinfer-cubin --index-url https://flashinfer.ai/whl     # this is the difference

pip install flashinfer-jit-cache --index-url https://flashinfer.ai/whl/cu129
# OR
pip install flashinfer-jit-cache --index-url https://flashinfer.ai/whl/cu130

See also #3808.

Gemma 4 and W4A16 on Blackwell RTX Pro and DGX Spark

Blackwell SM12x parts (RTX PRO 6000, DGX Spark / GB10) previously trailed Blackwell GB200 NVL72 (and B300 SM103) on attention shape and quantization coverage. This release closes the gap for Gemma 4 and weight-only-quantized inference: head_dim=512 attention for Gemma 4’s global layers now runs on SM120/121, the FMHAv2 prefill path gains head_dim=256/512 and sliding-window masking on SM120, the new mm_bf16_fp4 W4A16 GEMM is tuned for DGX Spark (completing W4A16 across both dense GEMM and MoE), and gated tanh-GELU brings Gemma 4 MoE onto the CUTLASS backend.

• #3576 feat(attention): head_dim=512 support for attention prefill & decode for Gemma 4 on SM120/121
• #3518 FMHAv2 on SM120 for head_dim 256/512 + sliding-window masks
• #3597 feat: Add BF16_FP4 GEMM with cuDNN and CuTe-DSL backends for SM120/121 for W4A16 workloads
• #3501 [feat] Add gated tanh-GELU (GeluTanh) activation to CUTLASS fused MoE (GEMMA 4)

GDN / gated delta rule: CuTe-DSL overhaul (Qwen3.5/3.6)

The gated-delta-rule kernels behind the Qwen3.5/3.6 family were rewritten from CUTLASS C++ to CuTe-DSL (#3491). The rewrite eliminates the C++ JIT compilation pain reported by customers and establishes the base for context-parallel delta-rule kernels — covering SM90 prefill and its context-parallel variant, SM120 prefill, and a ~20–25% GDN prefill speedup from mainloop efficiency work.

• #3477 feat: add sm90 delta rule dsl prefill
• #3481 feat: add sm90 cp delta rule dsl
• #3479 feat: add sm120 delta rule dsl prefill
• #3742 Improve GDN prefill perf by ~20-25% (mainloop efficiency)

GDN production-ready in vLLM

Two gaps blocked GDN serving in vLLM (#3602); both are now resolved. GDN kernels are compilation batch-size agnostic, so a single compiled cubin is reused across batch shapes instead of recompiling on every new batch size at inference time, and a new BF16 state recovery/decode kernel writes SSM state into preallocated space to supply the MTP-compatible spec-decode path vLLM needs.

• #3649 perf(gdn): make GDN kernels compilation batch-size agnostic (support dynamic batch shapes for vLLM integration)
• #3502 feat(gdn): BF16 state recovery/decode kernel with per-request K and f…

DeepSeek-class sparse MLA on Blackwell, FP8 KV on Hopper

New sparse-MLA paged-attention kernels extend the DeepSeek-V4 (d_qk=512) and DeepSeek-V3.2 / GLM-5.1 (d_qk=576) families onto SM120/121 through the existing flashinfer.mla APIs, with DSv4 coverage broadened to 8/16/32 head counts. On Hopper SM90, native FP8 KV cache support eliminates SGLang’s per-layer cast workaround, saving an HBM round-trip per layer on DeepSeek-V3/V4 while staying bit-identical to the BF16 path.

• #3395 feat(attention): add SM120 sparse MLA kernels
• #3545 Support smaller DSv4 sparse MLA head counts
• #3694 feat(attention): FP8 KV cache support for Hopper SM90 MLA