Release Highlights
Editorial highlights to inform users about important upgrades to their models, hardware, and serving stack.
v0.6.17
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.
v0.6.16
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
v0.6.15
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.
v0.6.14
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.