Sync with upstream v0.6.18 - #344
Merged
Merged
Conversation
…, fused quant+stage launch, persistent knob cache, and prequantized weight packs (flashinfer-ai#4079) <!-- .github/pull_request_template.md --> ## 📌 Description ### Summary Follow-up to flashinfer-ai#3980. This PR makes the MegaMoE (`moe_ep`) path consumable by serving engines (vLLM-style integration): full CUDA-graph capture/replay support, a fused single-launch quant+stage hot path, engine-friendly weight lifecycle (prequantized packs, post-preprocess source release), pooled symmetric-buffer workspaces shared across layers, and a persistent knob cache so production sessions resolve tuned knobs with a pure lookup. No in-engine autotuning. ### Performance Highlights All numbers measured on a **single node, 4x GB200 (EP=4)**. The built-in knob heuristic currently supports **GB200 only** (offline-tuned knob caches expected to work on any supported device). **Microbenchmark** — DeepSeek-V3 geometry (256 experts, top-8, hidden 7168, intermediate 2048), `e2e_pipelined` p50 µs, heuristic knobs, best variant (`nvfp4 + combine_nvfp4`) vs `deep_gemm_mega` baseline: | tok/rank | deep_gemm_mega | nvfp4 + combine_nvfp4 | Speedup | |---|---:|---:|---:| | 512 | 340.0 | 321.1 | 1.06x | | 1024 | 468.0 | 363.5 | 1.29x | | 2048 | 817.2 | 529.4 | 1.54x | | 4096 | 1473.5 | 862.7 | 1.71x | | 8192 | 2993.7 | 1677.3 | 1.78x | The sweep also covers real-model MoE geometries (DeepSeek V3 / V4-Flash / V4-Pro, Kimi K2.6, Qwen3.5-397B, gpt-oss-120b — the last enabled by the %64 alignment relaxation); the pattern holds everywhere: deep_gemm-parity at small token counts, 1.6–1.9x for fp4 combine-wire at large tokens on 7168-hidden shapes. Full sweep tables (all variants), accuracy, and benchmark methodology: [`kernel_src/cutedsl_megamoe/TUNING.md`](https://github.com/mhoqueanik/flashinfer-moe_ep/blob/fi-moe_ep-framework-integration/flashinfer/moe_ep/kernel_src/cutedsl_megamoe/TUNING.md#real-model-geometry-sweep-2026-07-21). **End-to-end (vLLM)** — early e2e benchmarks with this backend integrated into vLLM 0.25.1 (DeepSeek-V4-Flash, 4x GB200 TP4/EP4, CUDA graphs capturing all recurring step shapes including prefill chunks, per-role offline knob caches) show **~18% prefill and ~7% decode throughput gain over native vLLM** (i.e.: vLLMs built-in deep_gemm mega path): | Workload | native vLLM | fi MegaMoE (nvfp4) | Speedup | |---|---:|---:|---:| | Prefill, 8k-token chunks (tok/s) | 45,701 | 53,962 | 1.18x | | Decode, 1024-seq concurrency (output tok/s) | 21,049 | 22,614 | 1.07x | | GSM8K (200q, greedy) | 0.965 | 0.975 | — | Numbers reproduce within 1% at this branch tip; details and caveats (the decode number requires graph capture to cover prefill chunk shapes) are in TUNING.md's e2e section. The vLLM integration itself will land as a separate PR. **Accuracy** — per-variant microbench rel-L2 (`acc_loss_pct` vs fp32 dense-MoE reference) is unchanged from flashinfer-ai#3980; see the accuracy table in TUNING.md. On real-model distributions the nvfp4 path is at parity-or-better (GSM8K above), consuming the single-quant nvfp4 checkpoint directly via the prequantized weight-pack path. ### Key Changes **CUDA Graph Support (cutedsl mega path)** - Warmup contract + capture guards for graph-safe forward passes - Tail-mask memo made graph-capture-aware; records the actual staged count (fixes a view-slicing regression on capture-touched buffers) - 2-rank lockstep CUDA-graph replay test for the nvfp4 mega layer **Launch-Path Optimizations** - Fused single-launch quant+stage for the cutedsl mega path; fused staging for `deep_gemm` mega with a 16B-alignment fallback (bit-identical to the multi-kernel path, toggle: `FLASHINFER_MEGA_FUSED_STAGE`) - Zero-copy output + memoized tail mask + caller-owned ikr buffers to cut per-call launch count - Cached stream-aware launch thunk in the nvfp4 backend hot path **Engine Integration: Weights & Workspaces** - `MoEWeightPack` split into `Unquantized | Prequantized` variants - Source weight pack released after preprocess (fixes OOM at model load) - Pooled symmetric-buffer workspaces shared across mega layers **Tuning Infrastructure** - Persistent knob cache (`FLASHINFER_MOE_EP_KNOB_CACHE`): offline-tuned winners resolved by pure lookup before the built-in heuristic; populated via the new offline tuner CLI `python -m flashinfer.moe_ep.tune`. Per-role cache files (prefill-tuned / decode-tuned) are the validated deployment pattern. - Skew-aware tuning (`--skew`), schedule-axis sweep (`--sweep schedule`), `--live-tokens` **Shape Coverage & Guardrails** - cutedsl mega alignment relaxed to %64 (dg keeps %128) — enables gpt-oss-120b-class geometries - Warn-once import check for the CuTe-DSL 4.6.1 performance floor **Determinism Investigation (docs)** - Reported `fi_dg` cross-run nondeterminism traced to vLLM batch-formation timing (schedule diff, 1-tok vs 8-tok batch) — engine-level and backend-independent; closed as a FlashInfer issue with probe tests retained ### Directories Affected - `flashinfer/moe_ep/backends/mega/` — CUDA graph support, fused staging, launch caching - `flashinfer/moe_ep/kernel_src/cutedsl_megamoe/` — shim knob cache, TUNING.md - `flashinfer/moe_ep/core/`, `flashinfer/moe_ep/modes/` — workspace pooling, weight-pack lifecycle - `flashinfer/moe_ep/tune.py` — offline tuner CLI - `tests/moe_ep/` — CUDA graph (single- and multi-rank), fused stage, knob cache, weight-pack, workspace-pool, determinism-probe coverage ### Next Steps (WIP) - Generalized tuning: extend the knob heuristic / offline-tuned profiles beyond GB200 - Performance support for CuTe-DSL 4.5.2 (currently 34–54% slower than 4.6.1, which is treated as the perf floor with a warn-once import check) ## 🔍 Related Issues Follow-up to flashinfer-ai#3980 (MegaMoE kernel restructure). ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). New tests: `test_mega_cuda_graph.py`, `test_mega_cuda_graph_multirank.py`, `test_fused_quant_stage.py`, `test_knob_cache.py`, `test_weight_pack_union.py`, `test_workspace_pool.py`, `test_moe_ep_deep_gemm_skew_determinism.py`. Multi-rank tests run under `tests/moe_ep/run_tests.sh`. ## Reviewer Notes - The fused quant+stage path is bit-identical to the original multi-kernel path (`tests/moe_ep/test_fused_quant_stage.py`); `FLASHINFER_MEGA_FUSED_STAGE=0` reverts it for bisection. - The knob cache keeps autotuning strictly offline: `knobs=None` sessions do a pure cache lookup, falling back to the heuristic on miss — no measurement runs inside an engine process. <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added fused input staging for supported MoE quantization workflows, with an environment-variable toggle. * Added persistent autotuning cache support and an offline tuning command. * Added workspace sharing to reduce memory usage across compatible layers. * Added optional zero-copy output handling and CUDA Graph warmup support. * Added explicit pre-quantized and unquantized weight types. * **Bug Fixes** * Improved CUDA Graph capture safety, token tracking, alignment support, and workspace cleanup. * **Documentation** * Updated configuration and tuning guidance, including cache settings and performance information. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com>
## Summary - document FLASHINFER_CUTE_PREFILL_PERSISTENT in the CLAUDE.md environment-variable quick reference - clarify that 0 selects non-persistent scheduling and 1 selects persistent scheduling ## Testing - pre-commit run --files CLAUDE.md - git diff --check <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Documentation** * Documented the `FLASHINFER_CUTE_PREFILL_PERSISTENT` environment variable for controlling CuTe-DSL prefill scheduling. * Added explanations for the supported values: `0` for non-persistent scheduling and `1` for persistent scheduling. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
…lashinfer-ai#4175) ## Summary - add the missing `trtllm_fp8_per_tensor_scale_routed_moe` entry to `docs/api/fused_moe.rst` - fix the doc coverage regression reported by the doc checker ## Testing - not run (docs autosummary list change only) <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Documentation** * Added the routed FP8 per-tensor scale API to the TensorRT-LLM Fused MoE reference documentation. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
) ## Description Document the two `checkpointing_ssu` arguments still reported by the documentation checker in flashinfer-ai#4061: - `cu_seqlens`: packed variable-length boundaries, dtype/device requirements, and packed `x` shape - `max_seqlen`: required per-sequence upper bound, JIT specialization/ring-window role, and fixed-length exclusion The wording follows the wrapper validation and ring-buffer contract in the implementation. This is documentation-only and does not change runtime behavior. ## Validation - `git diff --check` - `pre-commit run --files flashinfer/mamba/checkpointing_ssu.py` - mypy passed - Ruff check/format passed - all remaining configured hooks passed Addresses the remaining `checkpointing_ssu` argument-consistency finding in flashinfer-ai#4061. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Documentation** * Clarified variable-length mode inputs for checkpointing. * Specified that packed sequence boundaries (`cu_seqlens`) use a 1D `int32` CUDA tensor of shape `(batch + 1,)`. * Updated required input shapes when sequence boundaries are provided, and clarified `max_seqlen` requirements (must be provided with boundaries, omitted otherwise). <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: hebo1221 <hebo1221@users.noreply.github.com>
<!-- .github/pull_request_template.md --> Fix ci failure issue introduced by incorrect testing tensor layout: tests/gemm/test_unified_gemm_fuzz.py Fix ci failure issue introduced by testing on unsupported cudnn version: tests/grouped_mm/test_grouped_mm_bf16.py tests/grouped_mm/test_grouped_mm_fp8.py ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). ## Reviewer Notes <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit - **Tests** - Improved NVFP4 quantization validation by aligning test inputs with the exact reference dequantization grid. - Updated grouped matrix multiplication tests to consistently detect supported cuDNN versions and capabilities. - Improved test skipping behavior for environments that lack required mixture-of-experts functionality. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Yanqin Zhai <yanqinz@nvidia.com>
…ointer packing (flashinfer-ai#3903) <!-- .github/pull_request_template.md --> ## 📌 Description Resolves item 14 of the SM121 support audit (flashinfer-ai#3170). The fused TRT-LLM allreduce module (`trtllm_comm`) refuses to build on SM12x because its JIT gate only lists SM90 and SM100, so those GPUs get no fused allreduce at all. The gate came from the original TRT-LLM port, not from a hardware limit. SM12x has everything these kernels use (verified on RTX 5080 and DGX Spark: thread-block clusters, distributed shared memory, programmatic dependent launch). Two real bugs sat behind the gate: - **Empty kernel on SM12x.** Separate from the JIT gate, the kernel header excludes SM12x with a compile-time check. Widening the JIT gate alone would build the lamport kernel (the low-latency path for small batches) with an empty body: small-batch decode calls would launch it, do nothing, and return garbage with no error. - **Wrong pointer layout in the host wrapper (all architectures).** The wrapper writes the lamport buffer pointers into a fixed-size array with one spacing between buffer sets, and the kernels read them back with another. The two layouts only agree at world size 16. At any other world size the kernel reads slots that were never written, then crashes or corrupts. No test ever caught this: the lamport kernel only runs at 16 tokens or fewer, and the existing tests use 64 and 128. Fixes: - The wrapper now writes the pointer array with the same spacing the kernels read. - The compile-time exclusion is removed, so the lamport kernel gets a real body on SM12x. - The JIT gate now includes major version 12, and AOT builds include the module for every arch it supports. - Launch cluster sizes are clamped to what the device reports through `cudaOccupancyMaxPotentialClusterSize`, instead of assuming a size per arch. SM90, SM100, SM120, and SM121 all report 8, so behavior is unchanged where the kernels already ran. - The test suites skip cleanly on unsupported machines instead of erroring. SM90 and SM100 compile to the same machine code as before; the compile-time check change does not affect them. ## 🔍 Related Issues flashinfer-ai#3170 (item 14, cross-cutting issue 5). ## 🧪 Tests On DGX Spark (GB10, SM121, CUDA 13, aarch64) and RTX 5080 (SM120, CUDA 13.3, x86_64): - JIT build through `gen_trtllm_comm_module` (`sm_121a` and `sm_120f`). - A functional test with two ranks sharing a single GPU over CUDA IPC. The lamport, oneshot, twoshot, and fusion paths all match fp32 references on both GPUs, including the small-batch lamport shapes the existing suites never reach. Before the pointer fix, those shapes crash or corrupt. - Both GPUs report max cluster size 8 for the real kernels and pass cluster-size-8 launches. - Cross-compilation for `sm_90a` and `sm_100a` is clean. - `pre-commit run` is clean. ## Reviewer Notes - The lamport dispatch thresholds (16 tokens or fewer, hidden size 256 or more) carry over from TRT-LLM's datacenter tuning. Whether lamport beats plain oneshot on SM12x needs a real multi-GPU measurement. - No API change. Test changes are only skip guards. 🤖 Generated with [Claude Code](https://claude.com/claude-code) <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Broadened TRT-LLM communication-kernel build support to additional GPU generations (SM90/100/100f/103/120/120f/121). * Added CUDA major version 12 handling for NVCC flag selection. * **Bug Fixes** * Added topology-limit validation for all-reduce IPC packing. * Corrected triple-buffer lamport IPC pointer packing stride. * Expanded architecture-guarded synchronization/launch-completion behavior. * Updated fusion and MoE cluster sizing to clamp using each device’s maximum supported cluster size. * **Tests** * Tests now run only when CUDA is available and compute capability is supported (9/10/12), with lamport buffers re-initialized per shape and expanded token coverage (including 8). <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
## 📌 Description Adds unified `UnpackedPrecomputed` routing support for `TrtllmFp4RoutedRunner`. Callers can provide separate contiguous routing tensors directly: - `topk_ids`: `[T, K]` `int32` global expert IDs - `topk_weights`: `[T, K]` BF16 and FP32 weights The runner forwards both tensors to the existing TRT-LLM FP4 Mode-3 launcher without constructing packed `(expert_id << 16) | bf16(weight)` IDs. Existing `PackedPrecomputed` and `FromLogits` behavior is unchanged. Coverage includes input validation, zero-copy forwarding, nonzero EP offsets, CUDA graph replay, and unified-fuzzer cases. ## 🚀 Pull Request Checklist ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing. SM100 results: ```bash pytest -v \ tests/moe/test_unified_moe.py::TestActivationPackValidation \ tests/moe/test_unified_moe.py::TestTrtllmFp4UnpackedContract \ tests/moe/test_unified_moe.py::TestTrtllmEPOffset 26 passed, 2 warnings ``` ```bash FLASHINFER_UMOE_FUZZ=1 \ FLASHINFER_UMOE_FUZZ_ONLY_SEED=900017 \ pytest -v tests/moe/test_unified_moe_fuzz.py::test_unified_moe_fuzz 2 passed, 4 warnings ``` ## Reviewer Notes This PR intentionally supports only the TRT-LLM FP4 unified runner. FP8 block, BF16, and MxInt4 launchers still use packed expert IDs and are out of scope. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit - **New Features** - Added support for unpacked precomputed routing in TRTLLM FP4 MoE execution. - Uses separate `int32` expert-id tensors and BF16/FP32 weight tensors (no repacking), preserving caller-provided tensors. - Enforces routing contracts for dtype, shape, contiguity, and unsupported-field rejection. - Extended autotuning and CUDA-graph execution to cover the new routing mode. - **Tests** - Expanded unit, fuzz, expert-parallel, autotuner, and CUDA-graph coverage for unpacked routing. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
…ai#4178) ## Why Kimi K3 MLA uses 96 global query heads with one KV head, TP-local head counts of 48/24/12/6, speculative query lengths up to 8, variable request-local query lengths, and context parallelism for long contexts. The existing Blackwell decode paths do not cover that combination efficiently or completely. ## Changes - Pack `(query_token, query_head)` rows continuously into CuTeDSL M128 tiles, so partial head groups can share a tile across tokens and only the final tile is padded. - Add compact variable-Q support to monolithic CuTeDSL with `[total_q, Hq, D]`, `cum_seq_lens_q`, and `max_q_len`. - Extend TRTLLM-GEN dense and sparse MLA to non-power-of-two `Hq <= 64` with padded Q8/Q16/Q32/Q64 tiles; `65 <= Hq <= 127` is routed to CuTeDSL. - Normalize split-KV to nonempty K partitions and parallelize underfilled D512 reduction across 1/2/4 output bands. - Classify dense versus causal work per M128 query tile so fully visible K tiles skip masking. - Add static cyclic decode context parallelism to monolithic BF16/FP16/FP8 CuTeDSL MLA. `cp_world` is compile-time, `cp_rank` is runtime, physical KV lengths remain rank-local, and `causal_seqlens_kv_global` supplies the global causal bound. The kernel returns rank-local output/LSE states for caller-side cross-rank merging. ## Correctness fixes - Close FP8 two-softmax named-barrier generations before a persistent 2-CTA cluster advances, fixing the upstream multi-work-item hang. - Bound partial-query TMA descriptors and predicate output/LSE tails. - Keep split count, workspace sizing, scheduler geometry, and autotuner keys consistent. - Handle DCP ranks or splits with no visible local keys as the neutral state: `O=0`, `LSE=-inf`. - Reject incompatible DCP combinations explicitly: DCP is monolithic fixed-Q only, requires `return_lse=True`, and cannot be combined with sinks or compact variable Q. ## Performance B200 CUDA-graph measurements for the packed-query and reducer changes showed: - up to 1.27x for flat query packing at B148/K8K; - 1.06x at B1 and 1.08x at B16 for H96/Sq8; - 9.7-13.0% end-to-end improvement from adaptive D2/D4 reduction; - no repeatable regression in a cold-L2 CUDA-graph sweep over `B={1,4,16,32,128}`, `K={1K,16K,128K}`, `Hq={12,24,48,96}`, `Sq={2,4,8}`, and BF16/FP8. DCP is statically specialized, so its disabled path does not add runtime DCP work. ## Validation Validated on B200 with CuTeDSL 4.5.0: - DCP matrix: 27 passed, covering BF16/FP16/FP8, `Hq={6,12,24,48,96,128}`, `Sq={1,4,8}`, `cp_world={1,2,4,8,16}`, heterogeneous/empty local KV, split-KV, PDL, and CUDA graphs. - Existing CuTeDSL MLA decode module: 678 passed. - Remaining MLA attention modules, including TRTLLM-GEN, DeepSeek MLA, paging, sparse MLA, XQA, and backend selection: 14,404 passed; 24,506 architecture-specific skips. - Full `pre-commit run --all-files`: passed, including clang-format, mypy, Ruff check, and Ruff format. ## Current limitations - CuTeDSL MLA requires SM100+, `Hq <= 128`, latent dimension 512, and RoPE dimension 64. - DCP requires monolithic fixed-Q and external cross-rank output/LSE reduction. - Runtime per-request variable split-KV is not qualified; automatic fixed split-KV is. - Partial final M128 rows remain compute lanes, although their memory accesses and stores are bounded or predicated. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added compact variable-length query support for MLA decode, including ragged query batches and inferred maximum query lengths. * Added static distributed context parallelism (DCP) support for monolithic MLA decoding. * Improved automatic backend selection for variable-query, DCP, and incompatible feature combinations. * Added support for non-power-of-two head configurations across MLA kernels. * **Bug Fixes** * Improved query masking, workspace sizing, output layouts, and handling of empty or inactive query rows. * Added clearer validation for unsupported or conflicting decoder options. * **Tests** * Expanded coverage for variable queries, DCP, CUDA graphs, FP8, backend routing, and boundary conditions. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Perkz Zheng <PerkzZheng@users.noreply.github.com> Co-authored-by: mingyangw <mingyangw@nvidia.com>
## Description Bump version to 0.6.16 for release. **Cut point:** `main` at `290c0918` (108 commits since `v0.6.15.post1`). ## Related Issues (Gated-by PRs) https://github.com/flashinfer-ai/flashinfer/issues?q=is%3Aopen+label%3Av0.6.16 + flashinfer-ai#4061 ## Reviewer Notes **API changes review** API changes since v0.6.15.post1, using `scripts/list_apis.sh` ```diff diff -u \ <(scripts/list_apis.sh -d -p --ref v0.6.15.post1) \ <(scripts/list_apis.sh -d -p) --- /dev/fd/63 2026-07-26 00:49:23.348331963 -0700 +++ /dev/fd/62 2026-07-26 00:49:23.348331963 -0700 @@ -300,6 +300,19 @@ x_out: torch.Tensor | None = None, mode: MixedCommMode | None = None, ) -> torch.Tensor: +[Global Functions] +@flashinfer_api +def quantized_all_reduce( + inp: torch.Tensor, + group: dist.ProcessGroup, + *, + scale_group: int = SCALE_GROUP_DEFAULT, + block_size: int | None = None, + num_warps: int | None = None, + max_num_blocks: int | None = None, + p2p_phase3: bool | None = None, + output: torch.Tensor | None = None, +) -> torch.Tensor: class MNNVLAllReduceFusionWorkspace(AllReduceFusionWorkspace): @flashinfer_api def checkpoint_prepare(self) -> None: @@ -314,6 +327,7 @@ ep_rank: int, ep_size: int, max_num_tokens: int, + eplb_stats_num_experts: int = 0, ): @@ -338,6 +352,10 @@ ep_size: int, top_k: int, num_experts: int, + enable_pdl: Optional[bool] = None, + eplb_local_stats: Optional[torch.Tensor] = None, + enable_rank_mask: bool = False, + active_rank_mask: Optional[torch.Tensor] = None, ): @@ -357,6 +375,12 @@ output_scales: Optional[torch.Tensor] = None, output_scalar_scale: float = 1.0, sf_layout: SfLayout = SfLayout.layout_linear, + output: Optional[torch.Tensor] = None, + *, + use_low_precision: bool = False, + enable_pdl: Optional[bool] = None, + enable_rank_mask: bool = False, + active_rank_mask: Optional[torch.Tensor] = None, ) -> torch.Tensor: @@ -367,6 +391,7 @@ metainfo: torch.Tensor, ep_rank: int, invalid_expert_id: int, + enable_pdl: Optional[bool] = None, ): @@ -376,6 +401,7 @@ max_num_tokens: int, total_dispatch_payload_size_per_token: int, combine_payload_size_per_token: int, + eplb_stats_num_experts: int = 0, ): class MoeAlltoAll: @@ -386,6 +412,7 @@ max_num_tokens: int, hidden_size: int, extra_payload_bytes_per_token: int = 0, + eplb_stats_num_experts: int = 0, ) -> int: @flashinfer_api def checkpoint_prepare(self) -> None: @@ -404,6 +431,8 @@ runtime_max_tokens_per_rank: int, invalid_token_expert_id: Optional[int] = None, expert_id_payload_index: Optional[int] = None, + eplb_local_stats: Optional[torch.Tensor] = None, + active_rank_mask: Optional[torch.Tensor] = None, ) -> list[torch.Tensor]: @flashinfer_api @@ -416,6 +445,10 @@ output_scales: Optional[torch.Tensor] = None, output_scalar_scale: float = 1.0, sf_layout: SfLayout = SfLayout.layout_linear, + output: Optional[torch.Tensor] = None, + *, + use_low_precision: bool = False, + active_rank_mask: Optional[torch.Tensor] = None, ) -> torch.Tensor: @flashinfer_api @@ -462,8 +495,8 @@ *, max_token_per_sequence: int, max_sequence_kv: int, - actual_seq_lens_q: torch.Tensor, - actual_seq_lens_kv: torch.Tensor, + actual_seq_lens_q: Optional[torch.Tensor] = None, + actual_seq_lens_kv: Optional[torch.Tensor] = None, block_tables: Optional[torch.Tensor] = None, causal: bool, return_lse: bool, @@ -475,14 +508,16 @@ batch_offsets_k: Optional[torch.Tensor] = None, batch_offsets_v: Optional[torch.Tensor] = None, batch_offsets_stats: Optional[torch.Tensor] = None, + batch_offsets_units: str = "elements", out: Optional[torch.Tensor] = None, lse: Optional[torch.Tensor] = None, is_cuda_graph_compatible: bool = False, backend: Optional[str] = None, o_data_type: Optional[torch.dtype] = None, ) -> tuple[torch.Tensor, Optional[torch.Tensor]]: + [Global Functions] -@flashinfer_api(trace=add_rmsnorm_fp4quant_trace) +@flashinfer_api(trace=add_rmsnorm_fp4quant_trace_dispatch) def add_rmsnorm_fp4quant( input: torch.Tensor, residual: torch.Tensor, @@ -497,6 +532,7 @@ output_both_sf_layouts: bool = False, block_scale_unswizzled: torch.Tensor | None = None, enable_pdl: bool | None = None, + y_out: torch.Tensor | None = None, ) -> Union[ Tuple[torch.Tensor, torch.Tensor], Tuple[torch.Tensor, torch.Tensor, torch.Tensor] ]: @@ -630,11 +666,12 @@ head_dim_qk, head_dim_vo=None, causal=True, - sm_scale=1.0, + sm_scale=None, q_data_type=torch.float16, kv_data_type=torch.float16, window_left: int = -1, variant: AttentionVariant | None = None, + window_right: int = -1, ) -> None: @flashinfer_api(trace=cute_dsl_batch_prefill_run_trace) @@ -644,7 +681,9 @@ k: torch.Tensor, v: torch.Tensor, out: Optional[torch.Tensor] = None, - ) -> torch.Tensor: + return_lse: bool = False, + lse: Optional[torch.Tensor] = None, + ): [Global Functions] @flashinfer_api(trace=rmsnorm_fp4quant_trace) def rmsnorm_fp4quant( @@ -882,7 +921,7 @@ bmm1_scale_log2: Optional[torch.Tensor], device: torch.device, ) -> Union[float, torch.Tensor]: -@flashinfer_api(trace=trtllm_batch_decode_trace) +@flashinfer_api(trace=trtllm_batch_decode_trace_dispatch) def trtllm_batch_decode_with_kv_cache( query: torch.Tensor, kv_cache: Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]], @@ -913,6 +952,7 @@ return_lse: bool = False, bmm1_scale_log2: Optional[torch.Tensor] = None, multi_ctas_kv_counter_buffer: Optional[torch.Tensor] = None, + enable_block_sparse_attention: bool = False, ) -> Union[ torch.Tensor, FP4Tensor, Tuple[Union[torch.Tensor, FP4Tensor], torch.Tensor] ]: @@ -960,6 +1000,7 @@ non_blocking: bool = True, fixed_split_size: Optional[int] = None, disable_split_kv: bool = False, + q_len_per_req: int = 1, global_override_indptr_cpu: Optional[torch.Tensor] = None, ) -> None: [Global Functions] @@ -972,6 +1013,8 @@ expert_ids: torch.Tensor, lora_indices: torch.Tensor, lora_stride: int, + *, + per_pair_input: bool = False, ) -> None: @@ -987,6 +1030,8 @@ slice_start_loc: torch.Tensor, output_slices: List[int], lora_stride: int, + *, + finalize: bool = True, ) -> None: @@ -1010,18 +1055,6 @@ output_dim: Optional[int] = None, ) -> torch.Tensor: [Global Functions] -@flashinfer_api -def interleave_moe_scales_for_sm90_mixed_gemm( - scales: torch.Tensor, - group_size: int = 32, -) -> torch.Tensor: - - -@flashinfer_api -def interleave_moe_weights_for_sm90_mixed_gemm( - weight: torch.Tensor, - quant_type: str = "fp4", -) -> torch.Tensor: @flashinfer_api(trace=cutlass_fused_moe_trace) def cutlass_fused_moe( input: torch.Tensor, @@ -1050,12 +1083,40 @@ use_mxfp8_act_scaling: bool = False, min_latency_mode: bool = False, use_packed_weights: bool = False, + use_wfp4afp8_humming: bool = False, tune_max_num_tokens: int = 8192, enable_pdl: Optional[bool] = None, activation_type: ActivationType = ActivationType.Swiglu, swizzled_input_sf: bool = True, use_fused_finalize: bool = True, + profile_ids: Optional[List[int]] = None, + workspace_buffer: Optional[torch.Tensor] = None, ) -> torch.Tensor: + + + max_num_tokens: int, + hidden_size: int, + intermediate_size: int, + num_experts_total: int, + top_k: int, + *, + x_dtype: torch.dtype, + weight_dtype: torch.dtype, + output_dtype: torch.dtype = torch.bfloat16, + activation_type: ActivationType = ActivationType.Swiglu, + tp_size: int = 1, + tp_rank: int = 0, + ep_size: int = 1, + ep_rank: int = 0, + min_latency_mode: bool = False, + use_deepseek_fp8_block_scale: bool = False, + use_w4_group_scaling: bool = False, + use_mxfp8_act_scaling: bool = False, + use_fused_finalize: bool = True, + use_packed_weights: bool = False, + use_wfp4afp8_humming: bool = False, + device: Optional[torch.device] = None, +) -> int: @flashinfer_api(trace=trtllm_bf16_moe_trace) def trtllm_bf16_moe( routing_logits: torch.Tensor, @@ -1433,6 +1494,7 @@ swiglu_alpha: float = DEFAULT_SWIGLU_ALPHA, swiglu_beta: float = DEFAULT_SWIGLU_BETA, swiglu_limit: float = DEFAULT_SWIGLU_LIMIT, + use_fused_finalize: bool = True, ): @flashinfer_api(trace=cute_dsl_moe_wrapper_run_trace) @@ -1450,6 +1512,8 @@ w2_weight_sf: torch.Tensor, w2_alpha: torch.Tensor, tactic: Optional[Tuple] = None, + *, + per_token_scale: Optional[torch.Tensor] = None, ) -> torch.Tensor: [Global Functions] @@ -1479,6 +1543,8 @@ swiglu_alpha: float = DEFAULT_SWIGLU_ALPHA, swiglu_beta: float = DEFAULT_SWIGLU_BETA, swiglu_limit: float = DEFAULT_SWIGLU_LIMIT, + *, + per_token_scale: Optional[torch.Tensor] = None, ) -> torch.Tensor: [Global Functions] @flashinfer_api(trace=fused_topk_deepseek_trace) @@ -1506,6 +1572,43 @@ ) -> Tuple[torch.Tensor, torch.Tensor]: [Global Functions] @flashinfer_api +def bgmv_moe_gemm1_lora_delta( + hidden_states: torch.Tensor, + w_ptr_a: torch.Tensor, + lora_stride_a: int, + w_ptr_b: torch.Tensor, + lora_stride_b: int, + topk_ids: torch.Tensor, + lora_ids: torch.Tensor, + rank: int, + intermediate_size: int, + *, + lora_dtype: torch.dtype = torch.bfloat16, + scale: float = 1.0, + out_dtype: torch.dtype = torch.bfloat16, +) -> torch.Tensor: + + +@flashinfer_api +def bgmv_moe_gemm2_lora_delta( + gemm1_activation_output: torch.Tensor, + expanded_idx_to_permuted_idx: torch.Tensor, + w_ptr_a: torch.Tensor, + lora_stride_a: int, + w_ptr_b: torch.Tensor, + lora_stride_b: int, + topk_ids: torch.Tensor, + topk_weights: torch.Tensor, + lora_ids: torch.Tensor, + rank: int, + hidden_size: int, + *, + lora_dtype: torch.dtype = torch.bfloat16, + scale: float = 1.0, + out_dtype: torch.dtype = torch.bfloat16, +) -> torch.Tensor: +[Global Functions] +@flashinfer_api def has_monomoe() -> bool: @flashinfer_api def get_scratchpad_size_bytes() -> int: @@ -1555,6 +1658,152 @@ interleave_up: bool = True, ) -> torch.Tensor: [Global Functions] +@flashinfer_api(trace=sm90_mixed_gemm_humming_weight_preprocess_trace_dispatch) +def preprocess_moe_weights_for_sm90_mixed_gemm_humming( + weight: torch.Tensor, + raw_scale: torch.Tensor, + max_range: int = 11, + *, + interleave: bool = True, +) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor]: + + +@flashinfer_api(trace=sm90_mixed_gemm_scale_interleave_trace) +def interleave_moe_scales_for_sm90_mixed_gemm( + scales: torch.Tensor, + group_size: int = 32, +) -> torch.Tensor: + + +@flashinfer_api(trace=sm90_mixed_gemm_weight_interleave_trace) +def interleave_moe_weights_for_sm90_mixed_gemm( + weight: torch.Tensor, + quant_type: str = "fp4", +) -> torch.Tensor: + + + w1_bf16: torch.Tensor, + w2_bf16: torch.Tensor, + *, + num_local_experts: int, + hidden_size: int, + intermediate_size: int, + device: Optional[torch.device] = None, + permute_cache: Optional[dict] = None, +) -> Dict[str, torch.Tensor]: + + + x: torch.Tensor, +) -> Tuple[torch.Tensor, torch.Tensor]: + + + weights: torch.Tensor, +) -> Tuple[torch.Tensor, torch.Tensor]: + + + w1_bf16: torch.Tensor, + w2_bf16: torch.Tensor, + *, + num_local_experts: int, + hidden_size: int, + intermediate_size: int, +) -> None: + + + w1_bf16: torch.Tensor, + w2_bf16: torch.Tensor, + *, + variant, + num_local_experts: int, + hidden_size: int, + intermediate_size: int, + device: Optional[torch.device] = None, +) -> Dict[str, torch.Tensor]: + + + hidden_states_bf16: torch.Tensor, + *, + variant, +) -> Tuple[torch.Tensor, torch.Tensor]: + + + scale: Union[float, torch.Tensor], *, name: str, device: torch.device +) -> torch.Tensor: + + + weights: torch.Tensor, +) -> Tuple[torch.Tensor, torch.Tensor]: + + + w1_bf16: torch.Tensor, + w2_bf16: torch.Tensor, + *, + hidden_states_scale_global: Union[float, torch.Tensor], + intermediate_scale_global: Union[float, torch.Tensor], + num_local_experts: int, + hidden_size: int, + intermediate_size: int, + device: Optional[torch.device] = None, +) -> Dict[str, torch.Tensor]: + + + hidden_states_bf16: torch.Tensor, + *, + hidden_states_scale_global: Union[float, torch.Tensor], +) -> Tuple[torch.Tensor, None]: + + + w1_bf16: torch.Tensor, + w2_bf16: torch.Tensor, + *, + num_local_experts: int, + hidden_size: int, + intermediate_size: int, + device: Optional[torch.device] = None, + permute_cache: Optional[dict] = None, +) -> Dict[str, torch.Tensor]: + + + x: torch.Tensor, group_size: int = 64, dim: int = -1 +) -> torch.Tensor: + + + w1_bf16: torch.Tensor, + w2_bf16: torch.Tensor, + *, + num_local_experts: int, + hidden_size: int, + intermediate_size: int, + device: Optional[torch.device] = None, +) -> Dict[str, torch.Tensor]: + + + weights: torch.Tensor, +) -> tuple[torch.Tensor, torch.Tensor]: + + + w1_bf16: torch.Tensor, + w2_bf16: torch.Tensor, + *, + num_local_experts: int, + hidden_size: int, + intermediate_size: int, + activation: str = "silu", + device: Optional[torch.device] = None, +) -> Dict[str, torch.Tensor]: + + + w1_fp4: torch.Tensor, + w1_blockscale: torch.Tensor, + w1_global_scale: torch.Tensor, + w2_fp4: torch.Tensor, + w2_blockscale: torch.Tensor, + w2_global_scale: torch.Tensor, + *, + activation: str, + source_format: str = "modelopt", +) -> Dict[str, torch.Tensor]: +[Global Functions] @flashinfer_api(trace=gated_delta_rule_decode_trace) def gated_delta_rule_decode_pretranspose( q: torch.Tensor, @@ -1628,6 +1877,7 @@ checkpoint_cu_starts: Optional[torch.Tensor] = None, checkpoint_every_n_tokens: int = 0, use_cp: Literal["auto"] | bool = "auto", + state_indices: Optional[torch.Tensor] = None, ) -> Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]]: [Global Functions] @flashinfer_api(trace=mm_bf16_trace) @@ -1753,7 +2003,7 @@ block_size: int = 16, use_nvfp4: bool = True, workspace_buffer: torch.Tensor = None, - tactic: int = -1, + tactic=-1, ): @@ -1796,8 +2046,6 @@ out: Optional[torch.Tensor] = None, backend: Literal["cudnn", "cublas", "cutlass", "auto"] = "cublas", ): - - A: torch.Tensor, B: torch.Tensor, A_scale: torch.Tensor, @@ -2077,6 +2325,52 @@ [Global Functions] +@flashinfer_api(trace=mm_nvfp4_svdquant_trace) +def mm_nvfp4_svdquant( + a: torch.Tensor, + b: torch.Tensor, + a_sf: torch.Tensor, + b_sf: torch.Tensor, + alpha: torch.Tensor, + d: torch.Tensor, + l1: torch.Tensor, + bias: Optional[torch.Tensor] = None, + out: Optional[torch.Tensor] = None, + backend: Literal["cutlass"] = "cutlass", + enable_pdl: Optional[bool] = None, +) -> torch.Tensor: + + + x: torch.Tensor, + pre_quant_scale: torch.Tensor, + global_scale: torch.Tensor, + enable_pdl: Optional[bool] = None, + backend: Literal["cutlass"] = "cutlass", +): +@flashinfer_api(trace=nvfp4_quantize_smooth_trace) +def nvfp4_quantize_smooth( + x: torch.Tensor, + pre_quant_scale: torch.Tensor, + global_scale: torch.Tensor, + enable_pdl: Optional[bool] = None, + backend: Literal["cutlass"] = "cutlass", +) -> Tuple[torch.Tensor, torch.Tensor]: + + +@flashinfer_api(trace=svdquant_linear_trace) +def svdquant_linear( + x: torch.Tensor, + weight_fp4: torch.Tensor, + weight_sf: torch.Tensor, + alpha: torch.Tensor, + pre_quant_scale: torch.Tensor, + l2t_smoothed: torch.Tensor, + l1_scaled: torch.Tensor, + global_scale: torch.Tensor, + bias: Optional[torch.Tensor] = None, + enable_pdl: Optional[bool] = None, +) -> torch.Tensor: +[Global Functions] @flashinfer_api(trace=grouped_gemm_nt_masked_trace) def grouped_gemm_nt_masked( lhs: Tuple[torch.Tensor, torch.Tensor], @@ -2219,6 +2513,20 @@ ) -> torch.Tensor: [Global Functions] @flashinfer_api +def moe_gemm_fp8_nt_groupwise( + a: torch.Tensor, + b: torch.Tensor, + a_scale: torch.Tensor, + b_scale: torch.Tensor, + m_indptr: torch.Tensor, + scale_granularity_mnk: Tuple[int, int, int] = (1, 128, 128), + scale_major_mode: Literal["MN"] = "MN", + backend: Literal["cute"] = "cute", + out: Optional[torch.Tensor] = None, + out_dtype: Optional[torch.dtype] = None, +) -> torch.Tensor: +[Global Functions] +@flashinfer_api def moe_gemm_mxfp8_nt_groupwise( a: torch.Tensor, b: torch.Tensor, @@ -2232,7 +2540,7 @@ out_dtype: Optional[torch.dtype] = None, ) -> torch.Tensor: [Global Functions] -@flashinfer_api +@flashinfer_api(trace=recurrent_kda_trace) def recurrent_kda( q: torch.Tensor, k: torch.Tensor, @@ -2252,16 +2560,18 @@ num_spec_tokens: Optional[int] = None, num_accepted_tokens: Optional[torch.Tensor] = None, output: Optional[torch.Tensor] = None, + initial_state_source: Optional[torch.Tensor] = None, + initial_state_indices: Optional[torch.Tensor] = None, + beta_is_logit: bool = False, ) -> tuple[torch.Tensor, Optional[torch.Tensor]]: [Global Functions] @flashinfer_api def checkpointing_ssu( state: torch.Tensor, - old_x: torch.Tensor, - old_B: torch.Tensor, - old_dt: torch.Tensor, - old_cumAdt: torch.Tensor, - cache_buf_idx: torch.Tensor, + x_cache: torch.Tensor, + B_cache: torch.Tensor, + dt_cache: torch.Tensor, + ring_start: torch.Tensor, prev_num_accepted_tokens: torch.Tensor, x: torch.Tensor, dt: torch.Tensor, @@ -2282,6 +2592,11 @@ cu_seqlens: Optional[torch.Tensor] = None, max_seqlen: Optional[int] = None, enable_pdl: bool = False, + cb_scaled: Optional[torch.Tensor] = None, + cumAdt_vec: Optional[torch.Tensor] = None, + cb_old: Optional[torch.Tensor] = None, + precompute_heads_per_cta: int = 0, + algorithm: str = "auto", ) -> torch.Tensor: [Global Functions] @flashinfer_api(trace=selective_state_update_trace) @@ -2643,6 +2958,105 @@ enable_pdl: bool | None = None, ) -> torch.Tensor: [Global Functions] +@flashinfer_api(trace=msa_proxy_score_trace) +def msa_proxy_score( + q: torch.Tensor, + k: torch.Tensor, + cu_seqlens_q: torch.Tensor, + cu_seqlens_k: Optional[torch.Tensor] = None, + *, + page_table: Optional[torch.Tensor] = None, + seqused_k: Optional[torch.Tensor] = None, + causal: bool = True, + max_seqlen_q: Optional[int] = None, + max_k_tiles: Optional[int] = None, + output: Optional[torch.Tensor] = None, + reduce_heads: bool = False, + q_offset=None, +) -> torch.Tensor: + +@flashinfer_api(trace=msa_proxy_score_fp4_trace) +def msa_proxy_score_fp4( + q_fp4: torch.Tensor, + k_fp4: torch.Tensor, + q_scale: torch.Tensor, + k_scale: torch.Tensor, + q_global_scale: float, + k_global_scale: float, + cu_seqlens_q: torch.Tensor, + cu_seqlens_k: Optional[torch.Tensor] = None, + *, + page_table: Optional[torch.Tensor] = None, + seqused_k: Optional[torch.Tensor] = None, + causal: bool = True, + max_seqlen_q: Optional[int] = None, + max_k_tiles: Optional[int] = None, + output: Optional[torch.Tensor] = None, + reduce_heads: bool = False, + q_offset=None, +) -> torch.Tensor: + +[Global Functions] +@flashinfer_api(trace=msa_sparse_decode_attention_trace) +def msa_sparse_decode_attention( + q: torch.Tensor, + k: torch.Tensor, + v: torch.Tensor, + q2k_indices: torch.Tensor, + *, + page_table: Optional[torch.Tensor] = None, + seqused_k: Optional[torch.Tensor] = None, + cu_seqlens_k: Optional[torch.Tensor] = None, + seqlen_q: int = 1, + causal: bool = True, + softmax_scale: Optional[float] = None, + return_softmax_lse: bool = False, + k_scale: Optional[torch.Tensor] = None, + v_scale: Optional[torch.Tensor] = None, + k_global_scale: Optional[float] = None, + v_global_scale: Optional[float] = None, + q_offset=None, + partial_dtype: Optional[torch.dtype] = None, + force_fused: Optional[bool] = None, +): +[Global Functions] +@flashinfer_api(trace=msa_sparse_attention_trace) +def msa_sparse_attention( + q: torch.Tensor, + k: torch.Tensor, + v: torch.Tensor, + q2k_indices: torch.Tensor, + cu_seqlens_q: torch.Tensor, + cu_seqlens_k: Optional[torch.Tensor] = None, + causal: bool = False, + softmax_scale: Optional[float] = None, + page_table: Optional[torch.Tensor] = None, + seqused_k: Optional[torch.Tensor] = None, + return_softmax_lse: bool = False, + k_scale: Optional[torch.Tensor] = None, + v_scale: Optional[torch.Tensor] = None, + k_global_scale: Optional[float] = None, + v_global_scale: Optional[float] = None, + q_offset=None, + return_temperature_lse: bool = False, + lse_temperature_scale: float = 1.0, +): + + + + + +[Global Functions] +@flashinfer_api +def msa_topk_select( + max_score: torch.Tensor, + topk: int, + num_valid_pages: Optional[int] = None, + output: Optional[torch.Tensor] = None, + force_begin_blocks: int = 0, + force_end_blocks: int = 0, +) -> torch.Tensor: +[Global Functions] @flashinfer_api(trace=rmsnorm_trace) def rmsnorm( input: torch.Tensor, @@ -2769,6 +3183,24 @@ beta: torch.Tensor, eps: float = 1e-6, ) -> torch.Tensor: + + +@flashinfer_api(trace=layernorm_quant_trace) +def layernorm_quant( + out: torch.Tensor, + input: torch.Tensor, + gemma: torch.Tensor, + beta: torch.Tensor, + scale: Union[float, torch.Tensor], + eps: float = 1e-6, +) -> None: + out: torch.Tensor, + input: torch.Tensor, + gemma: torch.Tensor, + beta: torch.Tensor, + scale: torch.Tensor, + eps: float = 1e-6, +) -> None: C, output_dtype, warps_m, ctas_per_row, bytes_per_ldg, kernel_cfg, occupancy ): @@ -3576,6 +4008,17 @@ skip_softmax_threshold_scale_factor: float = 0, ) -> Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]]: [Global Functions] +@flashinfer_api(trace=silu_and_mul_nvfp4_quantize_trace) +def silu_and_mul_nvfp4_quantize( + input: torch.Tensor, + global_scale: torch.Tensor, + sf_vec_size: int = 16, + is_sf_swizzled_layout: bool = True, + is_sf_8x4_layout: bool = False, + enable_pdl: Optional[bool] = None, +) -> Tuple[torch.Tensor, torch.Tensor]: + + @flashinfer_api(trace=fp4_quantize_trace) def fp4_quantize( input: torch.Tensor, @@ -3628,6 +4071,7 @@ sf_vec_size: int = 16, ufp8_type: int = 1, is_sf_swizzled_layout: bool = True, + is_sf_8x4_layout: bool = False, ) -> torch.Tensor: @@ -3662,11 +4106,16 @@ a: torch.Tensor, backend: str = "cuda", enable_pdl: Optional[bool] = None, + sfLayout: SfLayout = SfLayout.layout_128x4, ) -> Tuple[torch.Tensor, torch.Tensor]: @flashinfer_api -def mxfp4_dequantize(a_fp4, a_sf): +def mxfp4_dequantize( + a_fp4, + a_sf, + sfLayout: SfLayout = SfLayout.layout_128x4, +): @flashinfer_api @@ -3776,6 +4225,13 @@ input: torch.Tensor, global_scale: torch.Tensor, sf_layout: int = SF_LAYOUT_128x4, + enable_pdl: bool | None = None, +) -> Tuple[torch.Tensor, torch.Tensor]: + + + input: torch.Tensor, + global_scale: torch.Tensor, + sf_layout: int = SF_LAYOUT_128x4, enable_pdl: bool | None = None, ) -> Tuple[torch.Tensor, torch.Tensor]: ``` <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Chores** * Updated the application version from 0.6.15 to 0.6.16. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
<!-- .github/pull_request_template.md --> ## 📌 Description `test_xqa` intermittently fails its tolerance check on **DGX Spark (GB10, sm121)**. **Reproduced on DGX Spark and fixed by this PR.** Two races in the JIT `mha.cu` decode kernel, only exposed under GB10's tight few-SM scheduling (does not reproduce on sm120). ## Fixes 1. **mbarrier / multi-block-semaphore release ordering** not honored on GB10 → producer overwrites `smem.x` early / scratch not visible to the last CTA. Add three release fences, compile-gated to sm121 (`#if __CUDA_ARCH__ == 1210`), so other archs are unchanged. 2. **Multi-block output write missing the last-CTA gate**: every sub-sequence CTA wrote the output; with sliding-window an empty sub-sequence writes zeros and races the correct write. Gate on last-CTA (`ctaShouldWriteOut`), mirroring `mha_sm90.cu`. Arch-independent. ## Validation (DGX Spark, sm121) - Before: full `test_xqa` under `-P8` parallel stress reliably fails (down to ~8% elements passing). - After: 40× `-P8` full-suite runs → **zero failures**. Fixes flashinfer-ai#3658. ## 🔍 Related Issues flashinfer-ai#3658 ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). ## Reviewer Notes <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved synchronization during multi-block processing to prevent premature data overwrites. * Ensured partial results are fully visible before they are combined. * Prevented duplicate output writes when multiple processing blocks contribute to the same result. * Improved reliability and correctness for supported GPU workloads. <!-- end of auto-generated comment: release notes by coderabbit.ai --> Signed-off-by: Qidi Sang <200703406+qsang-nv@users.noreply.github.com>
…t dimension (flashinfer-ai#3882) ## 📌 Description The CI failure (test_bmm_mxfp8[...cudnn...-128-128-128-1], flaky cos-sim ~0.898 vs 0.9) was caused by the test misusing bmm_mxfp8: B quantized along n instead of the reduction dim k, and linear-layout scales fed to a cudnn graph that only reads F8_128x4-swizzled scales. That degraded accuracy from ~0.999 to mean ~0.92, leaving the unseeded test to cross the 0.9 threshold ~2% of the time at the smallest shape. ### Changes In `flashinfer/gemm/gemm_base.py`, we have added warnings when using unsupported scale layouts. Even though it was previously claimed that 2D non-swizzled / linear was supported, the kernels only ever read 1D 128x4-swizzled scales. Note that we have chosen to NOT throw on these paths, to avoid breaking some unforeseen use-case out there. Nevertheless, as the tests show, 2D non-swizzled / linear result in wrong output. | Backend | How scale is read | Swizzle | Linear path? | If a wrong (linear/2D) scale is passed | |---|---|---|---|---| | **cutlass SM100** | flat ptr + `Sm1xxBlkScaledConfig::tile_atom_to_shape_SF*(problem_shape)`, vec 32; scale strides never read | 128x4 fixed | **no** | warns only | | **cutlass SM120/121** | flat ptr + compile-time `Sm1xxBlockScaledConfig<32>` atom | 128x4 fixed | **no** | **hard reject** — Python req + C++ `ICHECK_EQ(ndim,1)` | | **cute-dsl (SM100)** | `make_ptr` + fixed 6D `BlockScaledBasicChunk` layout `((32,4),(32,4))` from `sf_m/sf_n/sf_k` | 128x4, vec 32 | **no** | **hard reject** — `ValueError` if `ndim!=1` | | **trtllm (SM100/103)** | raw ptr + `SfLayout` **enum** as cubin-selection key; no `Linear` cubin compiled | **A: 128x4 or 8x4; B: 128x4 (+shuffle)** | **no** | warns only | | **cuDNN** | 3D descriptor + hardcoded `tensor_reordering.F8_128x4` over a 1D swizzled buffer | 128x4 fixed | **no** | mm: reject (`ndim==1`, `use_8x4=False`); bmm: no ndim gate | In `examples/pytorch/flashinfer_modules.py`, we have corrected the example to use `bmm_mxfp8` by quantizing along the correct axis and dropping the superfluous and possibly mis-interpreted batch dimension from the scales. In `benchmarks/routines/gemm.py`, we remove the dead code for linear scales as they are never supported and passing correctly swizzled scales on the `auto` backend path. In `tests/gemm/test_bmm_mxfp8.py` and `tests/gemm/test_mm_mxfp8.py`, we remove the linear scale paths, fix the scale layout on the other paths and tighten the tolerances now that usage is correct. ## 🔍 Related Issues Example logs of failure in CI: ``` 2026-07-05T14:09:41.327317Z 01O tests/gemm/test_bmm_mxfp8.py ........................................... [ 7%] 2026-07-05T14:09:41.327318Z 01O ...............................F........................................ [ 19%] 2026-07-05T14:09:41.327319Z 01O .............................sssssssssssssssssssssssssssssssssssssssssss [ 32%] 2026-07-05T14:09:41.327319Z 01O ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 44%] 2026-07-05T14:09:41.327320Z 01O sssssssssssssssssssssssssssss........................................... [ 57%] 2026-07-05T14:09:41.327333Z 01O ........................................................................ [ 69%] 2026-07-05T14:09:41.327333Z 01O .............................sssssssssssssssssssssssssssssssssssssssssss [ 82%] 2026-07-05T14:09:41.327334Z 01O ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 94%] 2026-07-05T14:09:41.327334Z 01O sssssssssssssssssssssssssssss [100%] 2026-07-05T14:09:41.327335Z 01O 2026-07-05T14:09:41.327335Z 01O =================================== FAILURES =================================== 2026-07-05T14:09:41.327336Z 01O E AssertionError: Cosine similarity 0.8984 is too low (expected > 0.9) 2026-07-05T14:09:41.327337Z 01O assert tensor(0.8984, device='cuda:0', dtype=torch.bfloat16) > 0.9 2026-07-05T14:09:41.327337Z 01O ----------------------------- Captured stderr call ----------------------------- 2026-07-05T14:09:41.327338Z 01O 2026-07-05 05:42:58,019 - INFO - autotuner.py:651 - flashinfer.jit: [Autotuner]: Autotuning process starts ... 2026-07-05T14:09:41.327338Z 01O 2026-07-05 05:42:58,021 - INFO - autotuner.py:674 - flashinfer.jit: [Autotuner]: Autotuning process ends 2026-07-05T14:09:41.327339Z 01O /workspace/flashinfer/tests/gemm/test_bmm_mxfp8.py:74: AssertionError: Cosine similarity 0.8984 is too low (expected > 0.9) 2026-07-05T14:09:41.327340Z 01O =============================== warnings summary =============================== 2026-07-05T14:09:41.327341Z 01O tests/conftest.py:16 2026-07-05T14:09:41.327341Z 01O tests/conftest.py:16 2026-07-05T14:09:41.327341Z 01O /workspace/flashinfer/tests/conftest.py:16: DeprecationWarning: tcgen05.OperandMajorMode is deprecated, use cute.nvgpu.OperandMajorMode instead 2026-07-05T14:09:41.327342Z 01O from cutlass.cute.nvgpu.tcgen05 import OperandMajorMode 2026-07-05T14:09:41.327343Z 01O 2026-07-05T14:09:41.327343Z 01O -- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html 2026-07-05T14:09:41.327344Z 01O -- generated xml file: /tmp/junit/tests_gemm_test_bmm_mxfp8.py.687030424.xml --- 2026-07-05T14:09:41.327344Z 01O =========================== short test summary info ============================ 2026-07-05T14:09:41.327345Z 01O FAILED tests/gemm/test_bmm_mxfp8.py::test_bmm_mxfp8[True-cudnn-res_dtype0-False-input_dtype0-128-128-256-1] - AssertionError: Cosine similarity 0.8984 is too low (expected > 0.9) 2026-07-05T14:09:41.327346Z 01O assert tensor(0.8984, device='cuda:0', dtype=torch.bfloat16) > 0.9 2026-07-05T14:09:41.327346Z 01O ====== 1 failed, 287 passed, 288 skipped, 2 warnings in 105.05s (0:01:45) ====== 2026-07-05T14:09:41.327347Z 01O ❌ FAILED: tests/gemm/test_bmm_mxfp8.py (pytest exit code: 1) ``` ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [X] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [X] I have installed the hooks with `pre-commit install`. - [X] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. ## 🧪 Tests - [X] Tests have been added or updated as needed. - [X] All tests are passing (`unittest`, etc.). ## Reviewer Notes <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit ## Summary * **Bug Fixes** * Improved MXFP8 GEMM/BMM diagnostics with clearer warnings when operand or scale tensor layouts don’t match expected contiguity/swizzle/column-major requirements. * Updated MXFP8 GEMM/BMM trace checking and reference math for 1D 128x4-swizzled E8M0 scales (including decode/unswizzle behavior). * Relaxed certain MXFP8 backend eligibility checks to warn instead of failing on non-critical scale-layout mismatches. * **Documentation** * Refreshed MXFP8 MM/BMM documentation and the PyTorch example to match supported swizzled scale layouts. * **Tests** * Increased cosine-similarity thresholds and aligned trace tolerance expectations. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Signed-off-by: jdebache <jdebache@nvidia.com>
<!-- .github/pull_request_template.md --> ## 📌 Description CuTe DSL MoE fused finalize overlaps its output memset on an auxiliary CUDA stream. During `autotune(True)`, the selected tactic is executed once more after `choose_one` returns but before the tuning context exits. Keeping the auxiliary memset enabled for that replay can leave cross-stream work outstanding and deadlock a subsequent first-time CUDA module load during multi-rank serving startup. This PR: - threads the existing internal `use_async_memset` control through the CuTe DSL MoE functional and wrapper dispatch layers; - disables the auxiliary-stream memset only for the selected-tactic replay while `AutoTuner.is_tuning_mode` is true; - preserves async memset for each timed tactic profile, normal inference, and explicit-tactic execution; and - adds a contract test covering functional/wrapper dispatch both inside and outside tuning mode. This follows the existing tuning-aware dispatch pattern in `flashinfer/mla/_sparse_mla_sm120.py`, which branches on `AutoTuner.is_tuning_mode`. It also relies on the distinction documented in `flashinfer/autotuner/autotuner.py`: `is_tuning_mode` covers the final invocation after `choose_one`, while `is_in_profile_measurement` covers only an individual tactic's measurement window. ## 🔍 Related Issues Downstream integration: sgl-project/sglang#28354 ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). Focused CuTe DSL MoE contract and GPU autotune coverage on B200: ```bash FLASHINFER_DISABLE_VERSION_CHECK=1 \ FLASHINFER_CUDA_ARCH_LIST=10.0a \ CUDA_VISIBLE_DEVICES=0 \ FLASHINFER_NVFP4_4OVER6=1 \ FLASHINFER_NVFP4_4OVER6_E4M3_USE_256=1 \ FLASHINFER_NVFP4_4OVER6_ERR_MODE=MSE \ FLASHINFER_NVFP4_4OVER6_ERR_USE_FAST_MATH=1 \ python3 -m pytest -v -s \ tests/moe/test_cute_dsl_fused_moe.py::TestAutotuneReplayMemsetContract \ tests/moe/test_cute_dsl_fused_moe.py::TestCuteDslFusedMoeFunctional::test_with_autotune \ tests/moe/test_cute_dsl_fused_moe.py::TestCuteDslMoEWrapper::test_wrapper_with_autotune ``` Result: `7 passed`. The downstream 4-GPU Nemotron NVFP4-online + EAGLE startup was also validated with cold target and draft MoE autotune caches, fused finalize left at its default, and CUDA graphs enabled. All ranks completed FlashInfer autotuning, target verification graph capture, draft decode/extend graph capture, and a health generation request. ## Reviewer Notes - This changes no public API or kernel implementation. It reuses `_moe_core_impl`'s existing `use_async_memset` control. - Only the post-selection replay within `autotune(True)` becomes single-stream. Timed tactic measurements still profile the production async path. - Explicit-tactic dispatch remains unchanged. - No collective or all-reduce behavior is changed by this PR; that path was not required for the fix. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Performance** * Improved MoE auto-tuning behavior by disabling asynchronous memory clearing during timing runs. * Asynchronous memory clearing remains enabled during normal execution of the selected tactic. * **Tests** * Added coverage for both functional and wrapper APIs across tuning and normal execution modes. * Verified output shape and memory-clearing behavior during tactic replay. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
…er-ai#4183) ## Why NCCL-EP gained fault tolerance ([`nccl_ep.h` `enable_mask` + `ncclEpMask*`](https://github.com/NVIDIA/nccl/blob/master/contrib/nccl_ep/include/nccl_ep.h#L227-L232)): a peer that times out during dispatch/combine is masked and skipped instead of tripping a GPU `trap()`. NIXL-EP has the equivalent (`update/query/clean_mask_buffer`). **`flashinfer.moe_ep` exposed neither** — it never set `enable_mask`, never called a mask API, and gave callers no way to learn a rank had died. One slow or dead EP rank killed the whole job. This adds the FT surface over **both** transports. `moe_ep` is pure Python, so there is no CUDA/C++ change here. --- ## How the FT API is used FT is opt-in per Fleet via a **knob**, not new `FleetParams` fields — `FleetParams` is the frozen *sizing* dataclass, while optional transport features already live in the knob namespace (cf. `FleetAlgoKnobTopologyCapacity`). The runtime API lands on the `Fleet` ABC as **concrete raising defaults**, not `@abstractmethod`, so no existing or out-of-tree Fleet breaks. ### 1. Probe, then enable ```python from flashinfer.moe_ep import ( FleetAlgoKnobFaultTolerance, FleetAlgoKnobTopologyCapacity, MoEEpLayer, supports_fault_tolerance, ) # Needs more than the backend being built: nccl_ep also needs an nccl4py with # GroupConfig.enable_mask AND a libnccl_ep exporting ncclEpMask*. Never raises. assert supports_fault_tolerance("nccl_ep") knobs = [FleetAlgoKnobFaultTolerance(timeout_ms=5000)] # 0 = transport default # nixl only: size the capacity for the largest world you will ever reach knobs.append(FleetAlgoKnobTopologyCapacity(n=32)) layer = MoEEpLayer(bootstrap, fleet_params, weights, fleet_knobs=knobs, backend=SplitConfig(comm=NCCLEPConfig(), kernel=IdentityConfig())) ``` `FleetAlgoKnobFaultTolerance(enabled=True, timeout_ms=0, reconcile_timeout_s=30.0, coordinator_takeover_s=10.0)`. **LOW_LATENCY only** on both transports — `validate_fleet_params` rejects FT + HIGH_THROUGHPUT at construction, because nccl leaves the mask buffer NULL under HT and the mask APIs then *abort the process*. ### 2. Serve, poll, recover ```python fleet = layer._ensure_fleet() for step in ...: layer.forward(t) # Between iterations only — both transports read the mask LIVE from the # dispatch/combine kernels, so mutating it mid-collective is a race. if fleet.query_fault(): # free on nccl (pinned host flag) agreed = fleet.reconcile_active_mask() # store-collective, death-tolerant fleet.clear_faults(readmit=False) # re-arm detection; keep serving DEGRADED # ... later, if the peer comes back (collective over survivors, blocking): # fleet.clear_faults(readmit=True) ``` | Method | Collective? | Blocks host? | Stream-ordered? | |---|---|---|---| | `supports_fault_tolerance` | — | no | no | | `query_fault()` | local | nccl no / nixl yes (small D2H) | nccl no / nixl yes | | `query_active_mask(out=None)` | local | no | **yes** | | `set_active_mask(mask)` | local (must be applied identically everywhere) | no | **yes** | | `reconcile_active_mask()` | **store-collective**, tolerates dead ranks | yes (≤ timeout) | yes | | `clear_faults(readmit=False)` | local | no | no | | `clear_faults(readmit=True)` | **collective over survivors** | **yes** | yes | | `active_mask_epoch` | — | no | no | **Canonical mask: `int32[world_size]`, `1 = active`, CUDA tensor** — matching `ncclEpMaskQuery` and vLLM's `query_active_mask()` naming. ### 3. Rules that will bite you 1. **The steady state is read-only.** The transport discovers the fault and masks the peer; the application's job is to *notice* (`query_fault` / `query_active_mask`). `set_active_mask` is the exceptional reconciliation path, not a per-step call — most callers only reach it via `reconcile_active_mask()`. If you do write the mask, write it only between iterations (kernels read it live). 2. **All survivors must reconcile in the same iteration slot** — they must pass the same `active_mask_epoch`. 3. **`clear_faults(readmit=True)` and `update_topology()` are alternatives, not a sequence.** The former re-admits a merely-delayed rank on the *same* communicator; the latter destroys the group and builds a **new** `ncclComm_t`, which is the only way to add or replace a process. Re-admitting after a rebuild is meaningless; before one it is wasted work. (If you do both, `MaskClean` must come first — it needs a live handle on the current group.) 4. **No FT call during CUDA-graph capture** — but note **dispatch and combine themselves ARE safe to capture**: neither transport compacts the surviving ranks' layout, so they re-read the mask on every replay and a rank that fails later is still skipped. What must not be captured is a *decision about* fault state. `query_fault()` is a host read, not stream work, so it can't be captured at all — it returns the capture-time answer and freezes the branch taken on it into the graph forever. All four FT entry points raise on capture with a per-operation reason. *(Thanks @lrbison — an earlier draft of this said "stale offsets", which was wrong.)* 5. **A rank can be told it is dead.** `reconcile_active_mask()` raises `MoEEpRankEvictedError` when the survivors agreed *this* rank is gone. It can't apply that (a rank may not mask itself) and mustn't keep serving (peers stopped sending it tokens). 6. **Dropped tokens are not renormalized, and experts are not re-homed.** `y_degraded[t] == y_healthy[t] * Σ_{alive k} topk_weights[t][k]`. Both omissions are deliberate: implicit renormalization would add a kernel to every forward, hide a serving-quality event, and divide by zero when a token's whole top-k died; re-homing the dead rank's experts is an EPLB-style job that belongs to the framework. Keep serving on a partial mask with `reconcile_active_mask()` → `clear_faults(readmit=False)` — no `update_topology`, no new communicator. Runbook has the opt-in renormalization snippet. --- ## Call stacks ### Group/Buffer creation (where masking is switched on) ``` nccl_ep nixl_ep ------- ------- NcclEpFleet.__init__ NixlEpFleet.__init__ ├ _index_knobs -> self._ft ├ _index_knobs -> self._ft ├ validate_fleet_params(fault_tolerance=) ├ validate_fleet_params(fault_tolerance=) ├ _check_ft_supported() ├ nixl_ep.Buffer(..., timeout_ms=) [TypeError -> actionable] │ ├ dataclasses.fields(GroupConfig) ├ update_memory_buffers(cap, ...) │ │ -> needs "enable_mask" │ -> allocates mask_buffer[capacity], 0xFF-memset │ └ mask_ffi().available └ connect_ranks([0, world)) │ -> needs the 5 ncclEpMask* symbols -> unmasks [0, world); tail stays masked └ _build_group_config() ├ kwargs["enable_mask"] = True ├ kwargs["timeout_ns"] = timeout_ms * 1e6 (omitted when 0) └ nccl.ep.Group.create -> ncclEpCreateGroup -> allocates mask_buffer[nRanks] + pinned async-error flag ``` ### `query_fault()` ``` nccl_ep nixl_ep NcclEpFleet.query_fault NixlEpFleet.query_fault ├ reject_graph_capture("query_fault") ├ self.query_active_mask() (kernel + D2H) │ ^ NOT because it touches a stream -- │ ^ which itself rejects capture │ because it does NOT: a host read └ compare against self._ft_applied │ cannot be captured, so it would │ return the capture-time answer and (no host error flag on this transport, │ freeze the branch into the graph so the applied mask IS the state) └ mask_ffi().get_async_error(group) ├ group.get_async_error() [native first] └ ctypes ncclEpGetAsyncError(group.ptr, &o) -> libnccl_ep.so -> reads PINNED HOST flag: no stream, no sync, free ``` ### `query_active_mask()` ``` nccl_ep nixl_ep ├ _reject_graph_capture() ├ _reject_graph_capture() ├ _ft_bufs() -> device int32[world] ├ _ft_raw() -> device int32[CAPACITY] └ mask_ffi().mask_query(group, dev_ptr, str) ├ buffer.query_mask_buffer(raw) -> ncclEpMaskQuery │ -> nixl_ep_cpp -> kernel copy -> D2D copy of mask_buffer │ (asserts numel == max_num_ranks) -> ALREADY 1 = active, identity └ (raw[:world] == 0).to(int32) ^ NOT 1-raw: buffer is 0xFF-memset so untouched entries read back as -1 ``` ### `set_active_mask(mask)` ``` nccl_ep nixl_ep ├ _normalize_mask -> list[int], rejects ├ _normalize_mask (same guard) │ masking the local rank └ for each CHANGED rank r != self: ├ host.copy_(...) PINNED staging buffer buffer.update_mask_buffer(r, mask=(a==MASKED)) │ ^ pinned because Update is stream-ordered: -> atomicExch kernel, one launch EACH │ a pageable buffer mutated next call is -> so we push only the DIFF; a blind │ a use-after-write race range(world) loop would inject `world` └ mask_ffi().mask_update(group, host_ptr, str) launches into the steady state -> ncclEpMaskUpdate (HOST ptr, unlike Query) ``` ### `reconcile_active_mask()` — shared, transport-free ``` FaultToleranceMixin.reconcile_active_mask (both backends) ├ local = self.query_active_mask().cpu().tolist() ├ store = self._ft_store() -> resolve_rendezvous_store(subsystem="ft") └ reconcile_masks_via_store(store, rank, world, local, epoch=active_mask_epoch, ...) ├ 1. store.set("ft/gen{E}/local/{rank}", bytes(view)) ├ 2. poll ONLY ranks we still believe alive, until timeout_s ├ 3. elementwise-AND what arrived; mask believed-alive ranks that never reported ├ 4. coord = min(active); coord publishes via ATOMIC compare_set("ft/gen{E}/decision") │ everyone else adopts whatever that key holds <- kills split brain │ (coordinator itself dead -> mask it, re-elect, bounded by world_size) └ 5. _adopt(): raises MoEEpRankEvictedError if the decision masks US └ self.set_active_mask(agreed) -> backend-specific path above ``` Deliberately **not** a `torch.distributed` allreduce: that would hang on exactly the rank being masked out. With a store, a missing key *is* the death signal. ### `clear_faults()` ``` nccl_ep nixl_ep readmit=False: readmit=False: └ mask_ffi().error_clear -> ncclEpErrorClear └ return (no sticky flag to re-arm) readmit=True: readmit=True: ├ guard: a handle must exist ├ buffer.clean_mask_buffer() │ (MaskClean asserts on the LL staging │ -> zeroes ALL `capacity` entries, │ buffer -> would SIGABRT from C) │ marking the never-connected tail ACTIVE ├ warn under EXPERT_MAJOR ├ _ft_applied = [ACTIVE] * world │ (MaskClean computes reset offsets └ re-mask [world, capacity) <- fixes that bug │ assuming RANK_MAJOR) ├ mask_clean -> ncclEpMaskClean │ COLLECTIVE over survivors; internally cudaStreamSynchronize's └ error_clear -> ncclEpErrorClear ^ always paired: MaskClean does NOT clear the flag, which is exactly why `readmit` is a flag on one method rather than two a caller can mis-sequence ``` --- ## Three transport facts that shaped the code 1. **NIXL's polarity is "nonzero = masked", not "1 = masked".** The buffer is `0xFF`-memset at allocation (an untouched entry reads back as `-1`) and the kernels test `!= 0`. The normalization must be `(raw == 0)`; the obvious `1 - raw` yields `2` for never-connected capacity-tail ranks and silently poisons every downstream `sum()`/`bool()`. 2. **`clean_mask_buffer` zeroes all *capacity* entries**, marking the never-connected tail **active** — a live bug on any fleet sized above its world. 3. **`disconnect_ranks` is suffix-only**, so NIXL cannot evict a *middle* rank; masked-and-degraded is the terminal state there. Each has a regression test. ## nccl4py binding gap `nccl4py` binds `GroupConfig.enable_mask`/`timeout_ns` but its `Group` stops at `create`/`create_handle`/`destroy`/`.ptr` — the five mask functions are unreachable from Python. This adds a ctypes shim on `Group.ptr` that **tries a native `Group` method first**, so it retires itself with no call-site churn once those bindings land. Note the symbols live in **`libnccl_ep.so`, not `libnccl.so.2`**; we bind the process-global namespace, which is the only resolution guaranteed to be the same library the caller's group came from. ## Reconciliation, and a hole the tests found Each transport masks *locally* — NCCL-EP's header calls mask consistency "a framework-level concern" — so survivors can disagree, and disagreeing masks deadlock the next dispatch. The decision is published via a single atomic `compare_set`, which is what prevents a split brain when a straggler's key lands between two survivors' polls; there is a dedicated regression test for that timing. Writing those tests surfaced a case the design missed: a rank that is alive but that some peer already timed out on gets ANDed out of the group. It can't apply that decision and can't ignore it, so it now raises `MoEEpRankEvictedError`. ## Also fixed (incidental) - `validate_fleet_params` now rejects a nixl topology capacity **below** world size — it previously sailed through and went out of bounds inside the transport. - `update_topology` now rejects growing past the capacity, for the same reason. `test_update_topology_diffs_ranks` was exercising that invalid config (4→6 on capacity 4); it now passes `capacity=8`, keeping its intent. ## Testing `254 passed` in `tests/moe_ep/`; **166 FT-related tests pass**. The 6 failures in the full run are pre-existing and environmental (nvfp4/CuTeDSL ninja JIT) — confirmed failing identically on base `6258e522`. The protocol, shim and both backends' wiring are covered host-only (HashStore + threads; a fake ctypes library; fake transports), so they run in CI with no GPU. **Not yet run:** the 4-GPU tiers — `test_moe_ep_fault_tolerance_multirank.py` (stalls a middle rank and walks detect → reconcile → degrade → re-admit, asserting the degraded output *exactly* equals the surviving-weight scaling) and `smoke_ft_ep.py` (hard SIGTERM kill). This host has 1 GPU and no transport built. `bash tests/moe_ep/run_tests.sh ft` runs both. ## vLLM **No vLLM code here**, by design. vLLM already owns the contract (`support_fault_tolerance` / `query_active_mask` / `query_fault` + the per-step hook in `gpu_model_runner.py`); the FlashInfer-EP managers just hardcode `False`, and since the manager is transport-parameterized, one base-class change covers both LL backends. `docs/design_docs/vllm_moe_ep_integration.md` §8 records the target shape and the four things that commit must get right — including that vLLM's existing nixl/deepep `query_active_mask()` return *raw* buffers with the opposite polarity, and that `self._fleets` holds several fleets sharing one EP group so FT state must be hoisted to a primary fleet. ## Review follow-up (commit 9) @lrbison's review caught a real defect: `query_fault()` on nccl_ep had **no** capture guard, because I'd reasoned "host read, therefore capture-safe" — exactly backwards. Being a host read is *why* it can't be captured, so it silently returned the capture-time answer and froze the branch on it into the graph. NIXL's already raised (it goes via `query_active_mask`), so the same call behaved differently per backend. Both now raise, with a regression test each. The "stale offsets" rationale was also wrong and is replaced with per-operation reasons. ## Commits 1. `FleetAlgoKnobFaultTolerance` + Fleet FT API surface (no backend touched) 2. promote nixl_ep's store resolver to `core.bootstrap_utils` (pure move) 3. TCPStore-based active-mask reconciliation 4. ctypes shim for the `ncclEpMask*` API 5. wire nccl_ep `enable_mask`/`timeout_ns` + FT methods 6. wire nixl_ep `timeout_ms` + FT methods (carries the polarity/capacity/tail fixes) 7. multirank fault injection + FT smoke + `run_tests.sh ft` 8. docs + vLLM design note 9. `fix`: guard `query_fault` against graph capture + correct the capture rationale (review follow-up) 🤖 Generated with [Claude Code](https://claude.com/claude-code) <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added opt-in fault tolerance for MoE expert-parallel fleets, including timeouts and coordinator takeover. * Introduced public APIs to query faults, manage active-rank masks, reconcile fleet state, clear faults, and track mask epochs. * Added backend capability detection to ensure fault tolerance only activates when supported. * **Documentation** * Added design docs and an operational runbook describing recovery behavior, ordering rules, and transport-specific constraints. * Added a vLLM integration design note for future wiring into the fault-tolerance flow. * **Tests** * Added host-only unit tests, backend-specific mocks, NCCL ctypes-shim tests, multi-GPU end-to-end coverage, and FT smoke tests. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…shinfer-ai#4199) ## 📌 Description Follow-up to flashinfer-ai#4137, addressing [the review feedback](flashinfer-ai#4137 (comment)) that arrived after auto-merge. Thanks @qsang-nv for the detailed analysis. Issues: - Under programmatic dependent launch (PDL), the generic XQA kernel (`csrc/xqa/mha.cu`) read `q_cu_seq_lens` and the scale tensors before the acquire that makes a producer kernel's writes visible. `q_cu_seq_lens` drives the output row offset, so a stale read could write into the wrong request's rows. - On SM90 with fp8 KV cache, the small-batch layout of the Hopper XQA kernel (`csrc/xqa/mha_sm90.cu`, used when `q_seq_len * head_group_ratio <= 32`) hardcodes a causal draft mask, so a full draft mask silently returned causal results. Existing refchecks use random data and cannot tell the two modes apart within fp8 tolerance. - The XQA trace template declared a `q_cu_seq_lens` input that the trace reference silently dropped. Fixes: - Move the `q_cu_seq_lens` and scale-tensor loads below the PDL acquire in `mha.cu`. - Extend the SM90 fp8 fallback so small-batch speculative decode also runs on the generic kernel, matching the sliding-window and ragged-Q cases. Restoring the Hopper fast path is tracked in flashinfer-ai#4198. - Add a deterministic mask test: zero Q and K make each output row an exact mean of the visible V values, so any deviation from the requested mask fails loudly on every architecture. It runs two shapes, one that falls back to the generic kernel and one that stays on the Hopper kernel on SM90 fp8. - Remove `q_cu_seq_lens` from the trace template until the trace reference supports ragged Q. - Normalize the ragged-Q module key inside the module getter, and build a separate ragged variant only when an SM90 target is compiled; on other targets it is identical to the uniform module. - Document the SM90 fp8 fallback in the `xqa()`, `xqa_batch_decode_with_kv_cache`, and `trtllm_batch_decode_with_kv_cache` docstrings. ## 🔍 Related Issues flashinfer-ai#4198 (restore the Hopper fp8 fast path for speculative decode). Review thread: flashinfer-ai#4137. ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. ## 🧪 Tests - Added `test_xqa_batch_decode_mask_mode_deterministic` (causal/full mask, bf16/fp8 KV, two head-group shapes) with exact expected outputs. - On SM120 (RTX 5080): the new test, the ragged-Q and sliding-window suites (256 cases), and the trace suite (970 cases) pass. - SM90/SM100 are covered by CI. On SM90 fp8 the new test's `head_grp_size=16` shape exercises the Hopper kernel and the `head_grp_size=4` shape exercises the widened fallback. ## Reviewer Notes - Behavior change on SM90 with fp8 KV cache: small-batch speculative decode now runs on the generic kernel, including causal masks. The mask lives on the device, so dispatch cannot check its content without a sync. The Hopper kernel previously returned causal results regardless of the requested mask; flashinfer-ai#4198 restores that fast path. - The `SPEC_Q_SEQ_LEN` build specialization is no longer reachable at run time but is kept: the fast-path restoration in flashinfer-ai#4198 re-enables it. - Ragged workloads cannot be represented in trace dumps until the trace reference supports ragged Q; removing the inert input beats shipping a wrong reference. 🤖 Generated with [Claude Code](https://claude.com/claude-code) <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved speculative decoding handling for variable-length queries, including safer execution ordering for mask/sequence offsets. * Refined SM90 FP8 KV-cache kernel selection, extending the conditions that fall back to the generic kernel (ragged queries, attention sinks, sliding-window, and small head-group bound). * Improved compilation/caching behavior to better match supported speculative-decoding configurations. * **Documentation** * Updated the XQA documentation to clarify when the generic kernel is used for SM90 FP8 speculative decoding. * **Tests** * Added a deterministic test covering causal/full speculative-decoding masks for both BF16 and FP8 KV caches. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
…infer-ai#4039) ## 📌 Description Lets `msa_sparse_attention` and `msa_sparse_decode_attention` accept K/V views split from a paged KV cache that packs K and V in one `2 * head_dim` content dim per token. This is vLLM's KV-cache layout for MiniMax-M3 (vllm-project/vllm#44455), so it unblocks the vLLM integration of the MSA attend on SM120/SM121 (vllm-project/vllm#48994). Design: - The wrappers detect the split-view pattern and hand the kernel the whole packed cache; a `kv_packed` kernel variant reads K from the first half and V from the second half of each token's row. The contiguous path is unchanged. - Packed views are accepted for paged bf16/fp16/fp8 caches only; NVFP4 and non-paged inputs still require contiguous K/V, since vLLM's packed cache is paged and never NVFP4, so nothing produces those combinations. - The detection result is memoized so repeated calls with the same cache views skip it. - Callers probe the capability via the new `msa_ops.SUPPORTS_PACKED_KV` flag. - Also documents that `v_global_scale` applies to any KV dtype, which packed-cache callers rely on for the fp8 V descale. ### Performance Each cell is the median time of the same op over 30 iterations with random block selections on a packed cache, divided by the same measurement on a contiguous cache (1.00 = accepting the packed layout costs nothing). Rows cover decode and prefill shapes at the two KV-cache dtypes the packed path accepts, bf16 and fp8. topk is 16, the MiniMax-M3 default; the head shape (8 query / 2 KV heads) is kept small so the kernels stay short and any per-call overhead would show. B=1 cells vary by a few percent between runs. | shape (KV dtype, batch B, KV length S) | RTX 5080 | RTX PRO 6000 | GB10 | |---|---|---|---| | decode bf16 B=1 S=8192 | 1.01 | 0.98 | 0.80* | | decode fp8 B=1 S=8192 | 1.04 | 1.02 | n/a** | | decode bf16 B=8 S=8192 | 0.98 | 1.02 | 1.03 | | decode fp8 B=8 S=8192 | 1.02 | 0.99 | n/a** | | decode bf16 B=64 S=8192 | 0.98 | 1.01 | 0.99 | | decode fp8 B=64 S=8192 | 1.00 | 1.01 | n/a** | | decode bf16 B=256 S=8192 | 1.01 | 0.99 | 1.01 | | decode fp8 B=256 S=8192 | 1.00 | 1.02 | n/a** | | decode bf16 B=8 S=32768 | 1.01 | 1.03 | 1.09* | | decode fp8 B=8 S=32768 | 1.02 | 1.05 | n/a** | | prefill bf16 B=4 Q=512 S=8192 | 1.00 | 1.00 | 0.99 | | prefill fp8 B=4 Q=512 S=8192 | 1.00 | 1.00 | n/a** | \* Short GB10 runs vary a lot between repeats, likely because the board shifts clocks under its shared CPU/GPU power budget; read the starred cells as parity within that noise. \** fp8 KV kernels currently fail to compile on SM121 with cutlass-dsl 4.5.2; this reproduces without this PR. ## 🔍 Related Issues Follow-up to flashinfer-ai#3655 (MSA for SM120/SM121); enables the vLLM MiniMax-M3 integration in vllm-project/vllm#48994. ## 🚀 Pull Request Checklist ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). - New `tests/msa_ops/test_packed_kv.py`: decode and prefill on packed split views vs contiguous copies, across HND/NHD cache layouts, MQA and GQA, and bf16/fp8; plus the capability flag, the memoized detection, and rejection of strided views that are not packed splits. - RTX 5080 (two units) and RTX PRO 6000 (SM120): every packed case matches its contiguous reference bit-exactly, and the full `tests/msa_ops` suite passes. - GB10 (SM121): the bf16 cases pass bit-exactly; fp8 does not compile there (see the performance footnote). ## Reviewer Notes - Most of the diff is plumbing for the second compiled kernel variant, where mistakes fail loudly at compile or launch time. The exception is `_packed_kv_view` in `_common.py`: it reinterprets the caller's memory under new strides, so if it accepts a layout it should reject, the result is silently wrong attention output rather than an error. That function and its memoization are the parts to review closely. 🤖 Generated with [Claude Code](https://claude.com/claude-code) <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Extended sparse attention benchmarks with a `--kv_layout` option supporting `flat`, `paged`, and `packed`. * Added packed-KV cache view support for sparse prefill and decode. * Exported a capability flag so applications can detect packed-KV support. * **Bug Fixes** * Added validation for supported `kv_layout`/dtype combinations and for valid packed-KV view layouts. * Improved decode execution to correctly propagate packed-KV metadata. * **Tests** * Added SM120-121 tests covering packed-KV behavior, layout/dtype coverage, and invalid view rejection. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
…infer-ai#4159) Expose the existing TRTLLM-gen `MxFP4xMxFP8` (W4A8) and `MxFP4xBf16` (W4A16) kernels through the unified MoE API. ## 📌 Description This is PR 3 in the unified MoE quantization series for FP8 support: 1. flashinfer-ai#4026 — unified block-scale FP8 (merged) 2. flashinfer-ai#4091 — unified per-tensor FP8 (merged) 3. This PR — unified TRTLLM MXFP4×MXFP8/W4A8 and MXFP4×BF16/W4A16 ### Changes - Generalize `TrtllmFp4RoutedRunner` beyond NVFP4: - `QuantVariant.MXFP4`: `MxE2m1` weights × `MxE4m3` activations - `QuantVariant.W4A16`: `MxE2m1` weights × BF16 activations - Add variant-aware TRTLLM FP4 preparation: - MXFP4 weights with 32-element UE8M0 scales - MXFP8 activation preparation for W4A8 - BF16 activation preparation for W4A16 - Add shape, dtype, and scale-layout validation. - Add unified conformance and fuzzer coverage for packed and `FromLogits` routing. ### Support matrix - NVFP4 and MXFP4/W4A8: SM100, SM103 - W4A16: SM100 only (remains disabled on SM103, matching upstream xfail flashinfer-ai#1754) - SM120/121: separate b12x backends where available - SM107 is unsupported after flashinfer-ai#4171 ### Scope - No CUDA/C++ kernel changes; both modes already exist in the TRTLLM flat API. - CUTLASS W4A8 is out of scope. ## 🚀 Pull Request Checklist ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. ## 🧪 Tests SM100, CUDA 13 CI container: - `tests/moe/test_unified_moe_mxfp4.py` - 21 passed - Unified fuzzer, packed routing: - seeds `900017,900018` - 2 passed - Unified fuzzer, `FromLogits`: - seeds `900019,900020` - 2 passed
…eel installs (flashinfer-ai#3142) ## Summary Add CLI commands to install the matching flashinfer-jit-cache and/or flashinfer-cubin wheels from the FlashInfer wheel index. Also update documentation to reflect using `--index-url` for `flashinfer-cubin` instead of installing my pypi. ## What changed - detect the FlashInfer and CUDA versions and build the matching pip install command - document the new command in the CLI docs - add CLI tests for version/index selection and dry-run behavior - update docs to use `--index-url` for cubin wheel ## Why Packages is too large for PyPI, so users need a helper that can resolve the correct custom index-url and matching package version automatically. ## Validation - Click dry-run smoke test in verified the generated command resolves to with - was not available in the current environment, so the added CLI tests were not executed here Closes flashinfer-ai#3033 <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit - **New Features** - Added new CLI commands to install pre-built cubin and JIT-cache wheels. - Added `download-kernels` to install both kernel artifact sets together. - Improved CUDA-version handling with auto-detection, compatibility mapping, nightly installs, and `--dry-run`. - Added a hidden `download-jit-cache` alias for convenience. - **Documentation** - Updated README quickstart and artifact management CLI docs to use the new `flashinfer` subcommands. - Clarified what cubin downloads include and documented new install/download command behavior. - **Tests** - Added/expanded mocked CLI tests covering exact install invocations, dry-run behavior, CUDA fallback/validation, aliases, and failure handling. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: dierksen <dierksen@dierksen-spark.localdomain>
<!-- .github/pull_request_template.md --> ## 📌 Description Add SiTU activation support for TRTLLM-Gen MoE for MXFP4 × MXFP8 and NVFP4 × NVFP4 ## 🔍 Related Issues <!-- Link any related issues here --> ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). ## Reviewer Notes <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added SiTU gated activation support for FP4 MoE in routed and non-routed execution, with configurable per-expert `gemm1_alpha`, `gemm1_beta`, and `gemm1_clamp_limit`. * Extended FP4 MoE trace templates with an `activation_type` input to select SiTU. * **Bug Fixes** * Improved activation-type mapping/validation so SiTU uses the correct gated behavior. * Made MoE kernel/tile selection dtype-aware (activation + weight) and corrected per-token scaling eligibility during config generation. * **Tests** * Expanded MoE coverage for SiTU and optional per-expert parameters, including additional routed-logit and autotune regression checks. * **Chores** * Updated the remote BMM artifact reference and checksum. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Signed-off-by: Siyuan Fu <siyuanf@nvidia.com> Co-authored-by: Siyuan Fu <siyuanf@nvidia.com>
…groupwise GEMM (flashinfer-ai#4130) ## 📌 Description Adds **fused MoE FC1** (GEMM1 + SwiGLU in a single kernel) and a **warp-cooperative MoE scheduler** to the SM120 CuTe groupwise MoE GEMM entry. Fusion is exposed via an `is_gated` flag on the existing `moe_gemm_{fp8,mxfp8}_nt_groupwise` API. Follow-up to flashinfer-ai#3562 (MXFP8) and flashinfer-ai#3891 (FP8), sharing the same in-tree kernel package and JIT module. With `is_gated=True` the kernel reads the `2I` gate+up weight (up in the first `I` columns of N, gate in the second), applies `up * SiLU(gate)` in the epilogue, and writes the `I`-wide activated output — fusing away the `2I` intermediate global write+read and the separate activation launch. **Entry**: `flashinfer.grouped_mm.moe_gemm_{fp8,mxfp8}_nt_groupwise(a, b, a_scale, b_scale, m_indptr, scale_granularity_mnk, out=None, is_gated=False)`. `is_gated=True` requires even `b.size(1)` and halves output N to `I`. ## Benchmark 1: fused vs unfused - **fused** = `moe_gemm_*_nt_groupwise(..., is_gated=True)` — GEMM1 + SwiGLU in one kernel. - **unfused** = `moe_gemm_*_nt_groupwise(..., is_gated=False)` producing the `2I` gate+up, then a standalone Triton SwiGLU. - CUPTI kernel-level, warmup 20 + 50-iter median, exclusive serial. Quant / scale-transform / routing / output allocation are outside the timed window; the unfused window covers GEMM + SwiGLU back-to-back. - **Δ = fusion speedup = `(unfused_us / fused_us − 1) × 100%`** (positive = fused faster). FP8 and MXFP8 reported separately. FC1 shapes: **Qwen3.5-35B** (E=256, K=2048, 2I=1024) · **Qwen3-235B** (E=128, K=4096, 2I=3072) · **DeepSeek-V3** (E=256, K=7168, 2I=4096). DeepSeek-V3 tops out at MPE=512, the other two at 1024. ### RTX PRO 6000 Blackwell Server Edition (SM120a, 188 SM) #### Qwen3.5-35B | MPE | fp8 fused (µs) | fp8 unfused (µs) | fp8 Δ | mxfp8 fused (µs) | mxfp8 unfused (µs) | mxfp8 Δ | |---|---|---|---|---|---|---| | 1 | 415.7 | 417.2 | +0.4% | 425.9 | 428.9 | +0.7% | | 4 | 418.9 | 420.7 | +0.4% | 427.9 | 432.8 | +1.1% | | 8 | 422.2 | 424.6 | +0.6% | 430.2 | 437.0 | +1.6% | | 16 | 431.2 | 433.4 | +0.5% | 436.1 | 445.2 | +2.1% | | 32 | 444.1 | 447.8 | +0.8% | 443.8 | 459.3 | +3.5% | | 64 | 457.0 | 480.7 | +5.2% | 461.4 | 490.4 | +6.3% | | 128 | 486.1 | 544.5 | +12.0% | 494.4 | 555.2 | +12.3% | | 256 | 552.9 | 694.9 | +25.7% | 567.0 | 708.0 | +24.9% | | 1024 | 1568.5 | 2087.1 | +33.1% | 1555.9 | 2089.5 | +34.3% | #### Qwen3-235B | MPE | fp8 fused (µs) | fp8 unfused (µs) | fp8 Δ | mxfp8 fused (µs) | mxfp8 unfused (µs) | mxfp8 Δ | |---|---|---|---|---|---|---| | 1 | 1127.5 | 1117.2 | -0.9% | 1159.9 | 1155.0 | -0.4% | | 4 | 1133.3 | 1125.4 | -0.7% | 1164.7 | 1164.0 | -0.1% | | 8 | 1137.5 | 1133.3 | -0.4% | 1166.9 | 1173.7 | +0.6% | | 16 | 1148.1 | 1153.2 | +0.4% | 1180.5 | 1192.6 | +1.0% | | 32 | 1164.1 | 1177.4 | +1.2% | 1189.9 | 1213.6 | +2.0% | | 64 | 1194.3 | 1231.9 | +3.1% | 1220.9 | 1259.1 | +3.1% | | 128 | 1235.5 | 1338.1 | +8.3% | 1261.7 | 1365.8 | +8.3% | | 256 | 1329.5 | 1557.9 | +17.2% | 1369.7 | 1590.5 | +16.1% | | 1024 | 4859.1 | 5462.0 | +12.4% | 5021.2 | 5705.6 | +13.6% | #### DeepSeek-V3 | MPE | fp8 fused (µs) | fp8 unfused (µs) | fp8 Δ | mxfp8 fused (µs) | mxfp8 unfused (µs) | mxfp8 Δ | |---|---|---|---|---|---|---| | 1 | 4984.5 | 4969.1 | -0.3% | 5176.4 | 5179.1 | +0.1% | | 4 | 5001.5 | 4999.3 | -0.0% | 5195.6 | 5219.7 | +0.5% | | 8 | 5019.1 | 5031.8 | +0.3% | 5210.7 | 5254.8 | +0.8% | | 16 | 5079.4 | 5119.5 | +0.8% | 5252.0 | 5328.3 | +1.5% | | 32 | 5128.5 | 5211.5 | +1.6% | 5294.0 | 5413.6 | +2.3% | | 64 | 5337.5 | 5437.7 | +1.9% | 5386.0 | 5621.7 | +4.4% | | 128 | 5467.0 | 5788.3 | +5.9% | 5665.0 | 5974.9 | +5.5% | | 256 | 6817.4 | 7261.4 | +6.5% | 6493.9 | 7089.9 | +9.2% | | 512 | 13300.1 | 13823.2 | +3.9% | 12227.7 | 13102.6 | +7.2% | ### RTX PRO 5000 Blackwell (SM120a, 110 SM) #### Qwen3.5-35B | MPE | fp8 fused (µs) | fp8 unfused (µs) | fp8 Δ | mxfp8 fused (µs) | mxfp8 unfused (µs) | mxfp8 Δ | |---|---|---|---|---|---|---| | 1 | 455.3 | 465.1 | +2.2% | 471.4 | 477.5 | +1.3% | | 4 | 468.1 | 475.4 | +1.6% | 483.1 | 492.2 | +1.9% | | 8 | 475.6 | 481.2 | +1.2% | 496.0 | 499.4 | +0.7% | | 16 | 486.6 | 488.7 | +0.4% | 503.7 | 507.4 | +0.7% | | 32 | 497.3 | 503.9 | +1.3% | 512.8 | 521.7 | +1.7% | | 64 | 525.7 | 549.7 | +4.6% | 540.0 | 577.2 | +6.9% | | 128 | 573.1 | 633.3 | +10.5% | 600.7 | 652.4 | +8.6% | | 256 | 706.4 | 862.8 | +22.1% | 714.5 | 873.1 | +22.2% | | 1024 | 2790.3 | 3457.5 | +23.9% | 2817.7 | 3368.1 | +19.5% | #### Qwen3-235B | MPE | fp8 fused (µs) | fp8 unfused (µs) | fp8 Δ | mxfp8 fused (µs) | mxfp8 unfused (µs) | mxfp8 Δ | |---|---|---|---|---|---|---| | 1 | 1354.6 | 1366.7 | +0.9% | 1420.4 | 1406.5 | -1.0% | | 4 | 1373.0 | 1377.1 | +0.3% | 1438.3 | 1426.4 | -0.8% | | 8 | 1380.4 | 1381.6 | +0.1% | 1444.6 | 1436.1 | -0.6% | | 16 | 1399.2 | 1397.4 | -0.1% | 1454.8 | 1449.8 | -0.3% | | 32 | 1411.9 | 1410.7 | -0.1% | 1466.3 | 1468.7 | +0.2% | | 64 | 1462.5 | 1469.2 | +0.5% | 1492.5 | 1544.5 | +3.5% | | 128 | 1492.8 | 1639.3 | +9.8% | 1548.5 | 1694.6 | +9.4% | | 256 | 2017.4 | 2323.8 | +15.2% | 2116.8 | 2342.8 | +10.7% | | 1024 | 8278.5 | 9218.5 | +11.4% | 8306.9 | 9258.2 | +11.5% | #### DeepSeek-V3 | MPE | fp8 fused (µs) | fp8 unfused (µs) | fp8 Δ | mxfp8 fused (µs) | mxfp8 unfused (µs) | mxfp8 Δ | |---|---|---|---|---|---|---| | 1 | 6241.2 | 6235.1 | -0.1% | 6487.4 | 6451.4 | -0.6% | | 4 | 6254.4 | 6254.3 | -0.0% | 6499.1 | 6481.9 | -0.3% | | 8 | 6269.1 | 6290.3 | +0.3% | 6512.7 | 6510.1 | -0.0% | | 16 | 6337.1 | 6366.4 | +0.5% | 6555.2 | 6596.9 | +0.6% | | 32 | 6380.8 | 6442.6 | +1.0% | 6597.4 | 6682.0 | +1.3% | | 64 | 6885.3 | 6726.1 | -2.3% | 6701.6 | 6894.7 | +2.9% | | 128 | 7021.1 | 7308.6 | +4.1% | 7061.2 | 7408.3 | +4.9% | | 256 | 10298.8 | 10741.8 | +4.3% | 10245.9 | 10829.9 | +5.7% | | 512 | 20201.4 | 20958.2 | +3.7% | 19835.6 | 21038.7 | +6.1% | Fusion is neutral (occasionally slightly negative) at decode — small MPE, where FC1 grouped GEMM is weight-bandwidth / per-expert-tile-schedule bound and the `2I` intermediate round-trip is negligible against the GEMM floor — and grows with token count at prefill. Peak gain scales inversely with K: **Qwen3.5-35B (K=2048) > Qwen3-235B (K=4096) > DeepSeek-V3 (K=7168)** (up to +34% / +17% / +9% on the 6000; +24% / +15% / +6% on the 5000). FP8 and MXFP8 track closely. ## Benchmark 2: cooperative MoE scheduler (new `sm120_common/moe_scheduler.cuh`) The ZeroPadding token→tile map was previously resolved by a per-lane linear scan (base of flashinfer-ai#3891). This PR replaces it with a warp-cooperative scan (coalesced 32-group load + `shfl` prefix-sum + `ballot`, isomorphic to the CUTLASS SM90 group scheduler), selected per `GemmType` — non-ZeroPadding paths are SASS-unchanged. Decode-shaped microbenchmark on RTX PRO 6000 Blackwell Server Edition (M=1 token × topk=8, E=256, GranK=32; `contig` = padding-free upper bound), 4-arm NCU on one binary: | scheduler | GEMM1 (N=1024, K=2048) | GEMM2 (N=2048, K=512) | |---|---|---| | linear scan (prev, flashinfer-ai#3891) | 30.70 µs (+109% vs ideal) | 26.44 µs (+193%) | | **cooperative (this PR)** | **17.45 µs (+18.9%)** | **10.77 µs (+19.5%)** | | **speedup vs linear scan** | **+76%** | **+146%** | Scheduler global-load traffic drops **7.1×** (963,840 → 136,256 LSU sectors) and instructions 5×↓ (7.9M → 1.6M), bringing decode scheduler overhead from ~2–3× the padding-free ideal to within ~20%. At E≤64 both schedulers are within ~2–3% of ideal; large-M (prefill) and non-ZeroPadding shapes are unchanged (paired benchmark ≤0.42%, no regression). ## 🔍 Related Issues Follow-up to flashinfer-ai#3562 (MXFP8 MoE GEMM entry) and flashinfer-ai#3891 (FP8 MoE GEMM entry) — same in-tree kernel package. ## 🧪 Tests `tests/grouped_mm/test_cute_sm120_{fp8,mxfp8}.py` — adds gated (SwiGLU) correctness cells for both dtypes, against a per-expert bf16 GEMM + reference-SwiGLU baseline (FP8 `calc_diff < 2e-3`, MXFP8 `cos_sim > 0.99`). The `cute::Tensor` qualification in the blockscaling headers keeps existing non-gated coverage bit-identical. 🤖 Generated with [Claude Code](https://claude.com/claude-code) <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added optional gated MoE GEMM support for FP8 and MXFP8 workloads via `is_gated`. * Enables fused SwiGLU gate+up computation with reduced output width and stricter shape validation for packed gate/up inputs. * Improved SM120 MoE execution/tiling for gated workloads. * **Bug Fixes** * Tightened gated-mode dimension and alignment checks (including updated multiple-of-16 validation). * **Tests** * Added gated correctness coverage for FP8 and MXFP8 MoE GEMM, including relaxed numerical tolerance where needed. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…lashinfer-ai#4237) ## 📌 Description The routed dynamic-batch batched-GEMM kernels (`Bmm_*_dynB` with `-routeAct`) read one `int32` past the end of `ptrRouteMap`, from the last batch-dim CTA, on every launch. The read is speculative — its value is never consumed, so outputs are always bitwise correct — but it faults with `CUDA_ERROR_ILLEGAL_ADDRESS` whenever the allocation happens to end at a mapped-region boundary. That allocator-placement dependence is what makes it surface as *flaky* MoE autotune and inference crashes. Root cause is kernel-side: an off-by-one clamp in the hoisted load-task initializer (`WarpGrpThreadIdx` is clamped to the load-group size instead of size − 1), so off-group threads compute row `tileN` of a `tileN`-row tile, i.e. index `(ctaIdxY + 1) * tileN`. For the last CTA that is exactly one element past the shape documented in `KernelParamsDecl.h` (`[sum(divUpMul(N[bi], tileN) for bi in B)]`) — which is what we allocate. It has been reported to the kernel owners and is being fixed there. This PR is the integration-side workaround: allocate the route map with one extra element, so the already-shipped prebuilt cubins stay in bounds. The overrun is provably always exactly one `int32` from one CTA, so `+1` is sufficient by construction, not a heuristic. The pad slot's contents are irrelevant since the value is never used. Two allocation sites, both `permuted_idx_to_token_idx` (→ `routeMap`). Marked `WAR` + `TODO` so the `+1` can be dropped once regenerated cubins land. ## 🔍 Related Issues Likely explains flashinfer-ai#3530 and flashinfer-ai#3168 (and possibly flashinfer-ai#4012, flashinfer-ai#2776) — all report intermittent NVFP4 MoE autotune crashes, and flashinfer-ai#3168 additionally reports silent garbage output, which is the expected signature of an OOB read that usually lands on mapped memory. ## 🚀 Pull Request Checklist ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. ## 🧪 Tests - [ ] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). No test added — the failure is allocator-placement dependent and does not reproduce deterministically through the public API. Validation was done on `v0.6.15.post1` (B200, SM100, Kimi-K2.5 NVFP4 TP4): | Check | stock | route map +1 | |---|---|---| | `compute-sanitizer` memcheck, M=2 bucket × 4 ranks | 40 × `Invalid __global__ read of size 4` | **0 violations** (816 profiles/rank) | | Guard-page probe (route map placed at the tail of a VMM mapping) | faults on every launch | **clean** | | Full-model TP4 autotune crash loop | crashed by attempt 2 in 3/3 loops | **6/6 clean** | The guard-page probe is the load-bearing one: it removes all dependence on allocator placement, so it is deterministic in both directions. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved stability for fused mixture-of-experts routing operations by preventing out-of-bounds memory access in supported execution paths. * Applied the safeguard to both standard and FP4 block-scale routing workflows. <!-- end of auto-generated comment: release notes by coderabbit.ai --> Signed-off-by: Enwei Zhu <21126786+syuoni@users.noreply.github.com> Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
…oor (flashinfer-ai#4101) ## Summary `nvidia-cutlass-dsl==4.5.2` codegen has a regression that makes the cutedsl MegaMoE nvfp4 swap-AB fc12 kernel 34–54% slower than 4.6.1 at every token count (documented in `kernel_src/cutedsl_megamoe/TUNING.md`, "CuTe-DSL runtime sensitivity"). This mattered because vLLM 0.25.1 pins exactly 4.5.2, forcing integrations to carry a 4.6.1 force-upgrade plus its compat chain (quack, tvm-ffi, tilelang, vendored-kernel patches). This PR ports cutedsl_megamoe MR!27 (single file: `src/moe_nvfp4_swapab/kernel_fc12.py`): when the installed DSL is exactly 4.5.2, the MMA-consumer k-tile mainloop is peeled by one iteration (unconditional `try_wait` inside the loop, last tile issued after it). The gate is `cutlass.const_expr`-folded at trace time, so on any other version the generated kernel is byte-identical to before. ## Version support matrix (all measured, 4x GB200, default geometry) | DSL version | nvfp4 µs @1024/2048/8192 | status | |---|---|---| | 4.5.0 | — | fails at `cute.compile` (unsupported) | | 4.5.1 | — | unmeasured; shim warns (<4.5.2) | | 4.5.2 without WAR | 583.7 / 878.6 / 2579.5 | the regression | | **4.5.2 + this WAR** | **428.6 / 621.5 / 1896.4** | full parity | | 4.5.3 | 424.4 / 613.3 / 1933.3 | natively fast (upstream fix) | | 4.6.0 | 427.0 / 614.1 / 1922.0 | natively fast | | 4.6.1 | 428.5 / 625.6 / 1923.5 | reference | Support statement: **>= 4.5.2 at full performance; < 4.5.2 unsupported.** The regression existed only in 4.5.2 and was fixed upstream in 4.5.3, so the exact `== 4.5.2` gate never affects any other version. ## Validation (all on pinned 4.5.2, 2026-07-22) - **Correctness:** full `tests/moe_ep/run_tests.sh all` — all 8 sections PASS (unit, torch-oracle, split multirank, bf16/nvfp4/ht correctness, mega multirank, smoke). - **Microbenchmark:** full 5-variant × 7-token-point reference sweep reproduces the 4.6.1 reference within run noise at every cell; adopted as the new TUNING.md reference table. mxfp8 needs no WAR (never regressed — measured). - **vLLM 0.25.1 e2e (DeepSeek-V4-Flash, 4x GB200):** headline pair reproduces on vLLM's own 4.5.2 pin with no force-upgrade: prefill-8k **1.176x** native (53,623 vs 45,582 tok/s), decode-1k **1.068x** (34,263 vs 32,086 total tok/s) — within 0.7% of the 4.6.1-stack references. ## Changes - `kernel_src/cutedsl_megamoe/src/moe_nvfp4_swapab/kernel_fc12.py`: the MR!27 loop peel, `const_expr`-gated on `== 4.5.2`. - `kernel_src/cutedsl_megamoe/shim/__init__.py`: perf-floor warning now fires only below 4.5.2, with the measured support matrix in the docstring. - `kernel_src/cutedsl_megamoe/TUNING.md`: 2026-07-15 sensitivity section marked OBSOLETE (kept as record); reference tables re-measured on 4.5.2 and adopted; version-gap measurements and support statement added. AI-assisted (Claude Code): MR port, benchmark reruns, and docs. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit ## Summary by CodeRabbit - **Performance** - Improved NVFP4 MoE behavior for supported CuTeDSL versions, including corrected 4.5.2 handling. - Refreshed microbenchmark and vLLM end-to-end results, with clarified throughput metrics and updated guidance on valid “performance floor” claims. - **Compatibility** - Added an import-time warning for outdated CuTeDSL versions, with an environment variable to suppress it. - **Bug Fixes** - Fixed zero-token routing/staging to correctly reset masked outputs in both fused and non-fused paths. - **Documentation / Tests** - Expanded the benchmarking runbook and tuning instructions; added a GPU regression test for zero-token masking. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com>
(For example, this epsilon is in official DSV3.2 inference demo https://github.com/deepseek-ai/DeepSeek-V3.2-Exp, and it's currently applied when num_groups = 1 (GLM 5.x), but not this path. Add it for potential fix for instability when using f32 expert correction bias. Co-authored-by: Brayden Zhong <brayden@radixark.ai>
…nfer-ai#4138) 📌 Description Update the CuTe-DSL `nvfp4_quantize` implementation so `a_global_sf` accepts either a host-side float or a single-element device-side `torch.Tensor`. The implementation compiles separate kernel variants for host-scalar and device-tensor scales and includes the scale representation in the kernel cache key. The existing device-tensor path remains backward compatible and reads the scale inside the kernel without calling `.item()`. This avoids device-to-host synchronization and preserves CUDA Graph capture compatibility. This PR also updates the API documentation and adds tests for host-float parity, CUDA Graph capture, FP8 input, forced-TMA execution, and cache-key separation. 🔍 Related Issues Fixes flashinfer-ai#4112 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used my preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [[the pre-commit documentation](https://pre-commit.com/)](https://pre-commit.com/). 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing. Tested on an NVIDIA B200 with CUDA 12.9 and PyTorch 2.8.0+cu129: - All pre-commit hooks passed - 28 CuTe-DSL cache tests passed - 115 selected NVFP4 tests passed - 66 standard parity cases - 1 CUDA Graph regression case - 12 FP8 parity cases - 36 forced-TMA parity cases Two unrelated deprecation warnings were reported; there were no test failures. Reviewer Notes The device-tensor path intentionally remains supported for backward compatibility and CUDA Graph capture. Host-scalar and device-tensor inputs use separate compiled kernel variants and cache keys. `silu_and_mul_nvfp4_quantize_cute_dsl` is unchanged and remains outside the scope of this PR. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit - **New Features** - NVFP4 quantization now accepts the global scale as either a host-side floating-point value or a single-element device tensor. - Added support for both global-scale formats across CuTe-DSL linear, swizzled, and TMA quantization paths. - **Bug Fixes** - Improved compatibility for existing tensor-based scale inputs without unnecessary device transfers. - Ensured consistent results between host-float and device-tensor inputs. - Added CUDA Graph coverage for NVFP4 quantization to improve capture and replay reliability. <!-- end of auto-generated comment: release notes by coderabbit.ai --> Co-authored-by: Ka-Hyun Nam <knam@nvidia.com>
…le test (flashinfer-ai#4221) test_deep_gemm_mega_kernel_matches_torch_reference (added in flashinfer-ai#3980) fails under plain pytest with "ValueError: ... environment variable RANK expected, but not set": its WORLD_SIZE guard only skips when the variable is set and != 1, so without torchrun it falls through to dist.init_process_group(backend="nccl"), whose default env:// rendezvous requires torchrun's RANK/MASTER_* variables. This is only reachable on CI jobs that combine plain-pytest discovery of tests/moe_ep, a cu13 image with the EP stack (deep_gemm) installed, and a capability-10 GPU — i.e. the B300 cu130 unit-test job, where it currently errors on every run. Fix: when RANK is absent, self-bootstrap a 1-rank NCCL group via an explicit tcp://127.0.0.1:<free-port> init (rank=0, world_size=1) instead of env://. Deliberately avoids os.environ mutation so RANK/MASTER_* don't leak to later tests in the same pytest process; teardown is unchanged (conftest.pytest_sessionfinish destroys the group). The torchrun path is preserved and run_tests.sh still exercises it. Verified on GB200 (single GPU): both launch modes pass with rel_l2=0.0027 vs the torch oracle; the plain-pytest mode was confirmed with RANK/WORLD_SIZE/MASTER_ADDR/MASTER_PORT/LOCAL_RANK explicitly unset. AI-assisted (Claude Code). <!-- .github/pull_request_template.md --> ## 📌 Description <!-- What does this PR do? Briefly describe the changes and why they’re needed. --> ## 🔍 Related Issues <!-- Link any related issues here --> ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). ## Reviewer Notes <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
…en (flashinfer-ai#4108) ## 📌 Description Adds a native TRTLLM-GEN sparse MLA decode path for the shape with **no rotary tail**: `kv_lora_rank=512`, `qk_rope_head_dim=0`. For this shape the query carries no RoPE component, and both KV TMA descriptors address a **single 512-wide cache pool**, so the kernel needs the per-request active lengths to bound the sparse gather. Today `trtllm_batch_decode_with_kv_cache_mla` only accepts `deepseek_mla_dimensions` and `smaller_mla_dimensions`, so a `qk_rope_head_dim=0` request is rejected as an unsupported MLA dimension. This PR registers the new dimension set and threads an optional `sparse_mla_top_k_lens` tensor down to the launcher so the shape can be served natively. **Changes** - **`csrc/trtllm_fmha_kernel_launcher.cu`** — add an optional `sparse_mla_top_k_lens` argument to `trtllm_paged_attention_decode`. When the single-pool dynamic sparse MLA shape is detected (`sparse_mla_top_k_lens` present and MLA decode), pass the key cache as the sliding-window KV pool so the kernel reads the active per-token lengths. The launcher already rejects combining block-sparse attention with sparse MLA (`sparse_mla_top_k <= 0` check), so the two stay mutually exclusive. - **`flashinfer/mla/_core.py`** — register `nope_mla_dimensions` (`kv_lora_rank=512`, `qk_rope_head_dim=0`); require `sparse_mla_top_k > 0` and a `sparse_mla_top_k_lens` tensor for this shape; thread the autotune profiling length through the decode tuning config so different `top_k` values key distinct autotune configs; expose `sparse_mla_top_k_lens` on the public `trtllm_batch_decode_with_kv_cache_mla`. - **`flashinfer/decode.py`** — forward the new optional argument at the two existing kernel call sites. - **`flashinfer/trace/templates/attention.py`** — declare the optional `sparse_mla_top_k_lens` input on the sparse MLA decode trace template so the trace schema matches the kernel signature. The new argument is **optional and defaults to `None`**, so the `deepseek_mla_dimensions` / `smaller_mla_dimensions` decode paths are unchanged. `sparse_mla_top_k_lens` (one `int32` active length per query token) is supplied by the caller. ## 🔍 Related Issues None. ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). Static checks pass (`clang-format` / `ruff` / `mypy` via `pre-commit`). The new path has been exercised end-to-end in a downstream serving stack that computes `sparse_mla_top_k_lens` from the page table and drives this decode path. Happy to add a focused in-tree unit test for the `qk_rope_head_dim=0` dimension registration + argument threading — see Reviewer Notes. ## Reviewer Notes - **Backward compatibility**: `sparse_mla_top_k_lens` is optional and defaults to `None`; all existing callers and the two established MLA dimension sets keep their current behavior. - **Mutual exclusion**: block-sparse attention and sparse MLA are already mutually exclusive in the launcher (`sparse_mla_top_k <= 0` check), so the new single-pool path cannot be entered together with block-sparse. - **Autotune keying**: the profiling length is threaded through the decode tuning config and into the cache key, so a dense request (`len(inputs)==4`) and a sparse request (`len(inputs)==5`) resolve to distinct autotune configs rather than mis-keying. - I can add a unit test covering the dimension registration and the optional-argument threading if you'd like it in-tree — let me know the preferred test shape. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added support for native no-RoPE MLA decoding. * Added optional per-query sparse attention lengths for supported MLA decode workloads. * Added validation for sparse attention length tensor type, shape, device, and contiguity. * Integrated sparse MLA inputs with direct decoding and autotuning paths. * **Bug Fixes** * Improved handling of supported MLA head configurations during decode dispatch. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
<!-- .github/pull_request_template.md --> ## 📌 Description @HumansAnd SGLang needs `top_k_page_table_transform` to support independent score-window and page-table starts in its fused packed PAGED DSA path: - `row_starts` identifies the score window used for top-k selection. - `page_table_row_starts` identifies the page-table window used to translate the selected local indices. The current FlashInfer API applies `row_starts` to both operations, so it cannot represent this case. SGLang must therefore fall back to the SGL kernel even when `--dsa-topk-backend flashinfer` is selected. That produces correct indices, but it bypasses FlashInfer's deterministic, tie-break, and graph-safe fused top-k behavior for this path. This PR adds optional `page_table_row_starts` support. When it is omitted, FlashInfer continues to use `row_starts` for both operations, preserving the existing API behavior. The separate start is propagated through the Python and trace APIs, TVM FFI binding, radix and filtered implementations, deterministic post-sort, graph-safe dispatch, and the trivial `length <= k` path. Once SGLang adopts a FlashInfer release containing this API, it can remove the packed PAGED fallback and use FlashInfer fused top-k with the intended backend semantics. This PR does not change the SGLang call site. ### API Design For each output row `i`, define: ```text batch_i = row_to_batch[i] if row_to_batch is provided, else i score_start_i = row_starts[i] if row_starts is provided, else 0 page_start_i = page_table_row_starts[i] if page_table_row_starts is provided, else score_start_i length_i = lengths[i] ``` Top-k selection produces local offsets `local_idx[i, j]` in `[0, length_i)` by ranking: ```text input[i, score_start_i + local_idx[i, j]] ``` The page-table transform uses the same local offsets but an independent page-table origin: ```text output[i, j] = src_page_table[batch_i, page_start_i + local_idx[i, j]] ``` **`page_table_row_starts` affects only the page-table lookup used to produce output values.** It does not change the input score window, the selected local offsets, or the output tensor shape; it only changes which columns of `src_page_table` are gathered into the output. The argument responsibilities are therefore orthogonal: - `row_to_batch` selects only the row of `src_page_table`; multiple score rows may map to the same page-table row. - `row_starts` selects only the score-window origin. - `page_table_row_starts` selects only the page-table-window origin. It may be provided independently of `row_starts`; when omitted, it reuses `row_starts` for backward compatibility. If `length_i <= k`, all local offsets `0..length_i-1` are transformed and the remaining output positions are `-1`. Otherwise, exactly `k` local offsets are selected according to the existing deterministic and tie-break semantics. Each optional mapping/start tensor has shape `(num_rows,)`, dtype `int32`, and resides on the same CUDA device as `input`. Callers must satisfy `0 <= length_i`, `0 <= batch_i < src_page_table.shape[0]`, `0 <= score_start_i`, `score_start_i + length_i <= input.shape[1]`, `0 <= page_start_i`, and `page_start_i + length_i <= src_page_table.shape[1]`. This selection-plus-gather contract is not specific to SGLang: it represents any packed layout where score storage and lookup-table storage use different origins. Absolute starts are used instead of deltas so the API does not assume a relationship between the two windows or expose framework-specific metadata such as `cu_seqlens`. ## 🔍 Related Issues - SGLang DSA top-k backend integration: sgl-project/sglang#22851 - SGLang packed PAGED correctness fallback and backend-selection cleanup: sgl-project/sglang#32490 ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). Validation on an NVIDIA B200 using the editable source build and source JIT: ```text pre-commit run --all-files Passed python3 -m pytest -vv -s tests/utils/test_topk.py -k 'test_top_k_transform_with_row_starts' 48 passed, 1334 deselected, 2 warnings in 0.61s ``` The test extends the existing `test_top_k_transform_with_row_starts` Cartesian product across radix/filtered dispatch, graph-safe mode, deterministic mode, shared/separate starts, and both trivial and selected rows. ## Reviewer Notes Review focus is welcome on propagation through the deterministic post-sort and graph-safe filtered paths, where page-table translation occurs separately from score selection. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added optional `page_table_row_starts` support to the fused Top‑K page-table transform for independent per-row destination window offsets. * Updated tracing and reference implementations to model separate score-window (`row_starts`) and destination page-table-window (`page_table_row_starts`) offsets. * **Bug Fixes** * Corrected page-table addressing when the score and destination windows start at different offsets. * **Tests** * Expanded Top‑K transform coverage to include deterministic mode and separate `page_table_row_starts`, with additional trace-based reference checks. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
…-ai#3910) ## Summary Make the existing intranode `MixedCommHandler` all-gather and reduce-scatter paths checkpointable without adding a new collective kernel. - Add one `_vm_mapped` boolean and idempotent `checkpoint_prepare()` / `checkpoint_restore(comm_backend)` methods. - Reuse the existing VMM startup structure for initial mapping and restore. - During prepare, synchronize through the retained VMM `CommBackend` and release UC/MC physical, imported, mapped, and bound resources while retaining all CUDA VA reservations and the GPU peer-pointer table captured by CUDA graphs. - During restore, validate a fresh local `CommBackend`, recreate backing resources at the exact retained addresses, reset mixed-comm protocol memory, collectively fence, and retain that backend for the next prepare or attached shutdown. - Clear the retained backend after prepare so its caller-owned process group may be destroyed while the workspace is detached. - Use the same unmap/release and permanent-address-free helpers in explicit shutdown, including shutdown from an already detached state. This PR deliberately does **not** add a standalone symmetric all-gather kernel. It makes the existing mixed-comm all-gather and reduce-scatter kernels checkpointable. ## Lifecycle contract The public API matches flashinfer-ai#3727/flashinfer-ai#3745: ```python handler.checkpoint_prepare() handler.checkpoint_restore(fresh_comm_backend) ``` The current backend must remain valid through `checkpoint_prepare()`. After prepare returns, MixedComm no longer retains it and the caller may destroy its underlying process group. The fresh backend passed to restore is retained until the next prepare or attached shutdown. MixedComm does not own or destroy externally supplied process groups. No Python launch or CUDA graph replay guard is added. The caller/checkpoint orchestrator is responsible for not executing mixed-comm operations while the workspace is detached. ## Scope Checkpoint prepare/restore currently supports intranode mixed-comm workspaces only. Multi-node checkpoint calls fail during preflight without mutating the handler. Existing multi-node/NVSHMEM execution and normal shutdown remain unchanged. `CommBackend` is control-plane rendezvous only; it does not add multi-node NVSHMEM/RDMA checkpoint support. The implementation follows the earlier checkpoint lifecycle pattern: - no additional workspace wrapper; - no attachment-state enum; - no terminal/partial-recovery state machine; - fail-fast resource transitions; - successful prepare/restore idempotence derived from `_vm_mapped`; - upstream-style unannotated dynamic CUDA resource fields. The branch is one commit directly on the PR base. Final diff: ```text flashinfer/comm/mixed_comm.py | 221 lines changed tests/comm/test_mixed_comm_checkpoint.py | 215 lines added 2 files changed, 374 insertions(+), 62 deletions(-) ``` ## Validation ### Static and local - Full pre-commit: pass, including `mypy --all-files`, Ruff lint, and Ruff format. - Python compilation and `git diff --check`: pass. - Focused checkpoint test: exactly 1 test collected; expected hardware skip on SM89 because CUDA multicast is unavailable. - Existing mixed-comm suite: exactly 6 tests collected; expected SM89 capability skips. - Independent implementation review: approved. ### B200 CUDA process checkpoint/restore Validated exact commit `1c2ead3272289c2688be8d810cc019734dcc5149` on two NVIDIA B200 GPUs in `nscale-dev` using NVIDIA's official `cuda-checkpoint` utility. Before checkpoint, both ranks passed eager execution and CUDA graph capture/replay for all four combinations: ```text ALLGATHER x FUSED_OPT_WAITS_UC ALLGATHER x FUSED_OPT_WAITS_MC REDUCESCATTER x FUSED_OPT_WAITS_UC REDUCESCATTER x FUSED_OPT_WAITS_MC ``` Each rank called `checkpoint_prepare()` twice, verified it was detached, verified the backend reference was cleared, and verified exact UC/MC/pointer-table address equality. No operation was launched or replayed while detached. An external controller drove both rank PIDs through: ```text running -> locked -> checkpointed -> locked -> running ``` While both were `checkpointed`, `nvidia-smi --query-compute-apps` returned no compute processes. After CUDA restore/unlock, each rank: 1. created a fresh Gloo process group and `TorchDistBackend`; 2. called `checkpoint_restore()` twice; 3. verified the fresh backend was retained and all addresses matched exactly; 4. replayed the original four captured graphs with changed inputs and exact AG/RS results; 5. ran a second prepare/restore lifecycle using the retained fresh backend; 6. replayed all four original graphs again with another set of changed inputs; 7. completed attached shutdown while the fresh group remained valid; 8. let the caller destroy the fresh process group afterward. Both worker return codes were zero. The successful E2E runtime was approximately 41 seconds; the clean optimized SM100a JIT build took approximately 361 seconds. The temporary test pod was deleted after evidence collection. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added checkpoint restore for mixed communication with reusable GPU virtual-memory mappings. * Extended checkpoint restore to accept a communication backend. * Switched setup/restore/shutdown synchronization to backend-mediated coordination. * **Bug Fixes** * Improved checkpoint prepare/restore correctness, including repeated checkpoint cycles. * Preserved GPU address reservations across prepare and restore, ensuring stable teardown. * **Tests** * Added distributed CUDA tests covering checkpoint/restore, CUDA graph capture and replay, and shutdown. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
…#4280) <!-- .github/pull_request_template.md --> ## 📌 Description This PR relands SM 107 support to main branch (reverted in flashinfer-ai#4171) as well as some other release fixes. #### Cherry Picks - flashinfer-ai#4191 - flashinfer-ai#4189 - flashinfer-ai#4200 - flashinfer-ai#4215 - flashinfer-ai#4225 - flashinfer-ai#4230 - flashinfer-ai#4235 - flashinfer-ai#4226 - flashinfer-ai#4257 - flashinfer-ai#4258 - flashinfer-ai#4261 #### Other Changes - Rubin guards from flashinfer-ai#4252's conflict resolution (`TLLM_RUBIN_FEATURES`: SiTuGlu static_asserts + tile-192 advertisement, compiled out for the Rubin BMM pin) - Test-contract update: `test_unified_moe.py` arch assertions written post-revert (flashinfer-ai#4159) flipped to the restored contract (FP4/BF16 claim 107; FP8 stays 100/103) <!-- What does this PR do? Briefly describe the changes and why they’re needed. --> ## 🔍 Related Issues <!-- Link any related issues here --> flashinfer-ai#4107, flashinfer-ai#4164, reverts flashinfer-ai#4171 ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [ ] Tests have been added or updated as needed. - [ ] All tests are passing (`unittest`, etc.). ## Reviewer Notes <!-- Optional: anything you'd like reviewers to focus on, concerns, etc. --> <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **New Features** * Added support for Rubin/SM107 GPUs across GEMM, MoE, attention, quantization, sampling, and DeepGEMM workflows. * Added architecture-aware kernel selection, memory sizing, compilation, and artifact handling. * **Bug Fixes** * Improved validation and error messages for incompatible GPU architectures and invalid kernel configurations. * Clearly rejects unsupported NVFP4 KV-cache operations on SM107. * **Documentation** * Updated installation guidance with the SM107 architecture target. * **Tests** * Expanded architecture coverage and compatibility checks across GPU test suites. <!-- end of auto-generated comment: release notes by coderabbit.ai --> --------- Co-authored-by: Vinnie6167 <Vinnie6167@users.noreply.github.com> Co-authored-by: Ka-Hyun Nam <knam@nvidia.com> Co-authored-by: Alex Yang <aleyang@nvidia.com> Co-authored-by: Claude Opus 5 <noreply@anthropic.com> Co-authored-by: Jimmy Zhou <79552142+jimmyzho@users.noreply.github.com>
## 📌 Description Update the SM103 `mm_fp4` tests to match the current in-memory `AutoTuner.profiling_cache` schema. PR flashinfer-ai#4004 changed each cache value from `(runner_id, tactic, profile)` to `(tactic, profile)`, because the runner is identified by `ProfilingCacheKey`. The tests added by PR flashinfer-ai#4063 still unpacked the legacy three-item value and therefore raised: ```text ValueError: not enough values to unpack (expected 3, got 2) ``` This change: - unpacks `(tactic, profile)` in both SM103 test cases; - continues to validate the selected runner through `key.runner_class_name`; - preserves the native K768 and generic K128/K256 tactic-family assertions. This is a test-only change; no GEMM, dispatch, or autotuner production code is modified. ## 🔍 Related Issues - Follow-up to flashinfer-ai#4063 - Cache schema change: flashinfer-ai#4004 ## 🚀 Pull Request Checklist ### ✅ Pre-commit Checks - [ ] I have installed `pre-commit` by running `pip install pre-commit` (or used my preferred method). - [ ] I have installed the hooks with `pre-commit install`. - [ ] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. Targeted checks completed: - `ruff check tests/gemm/test_mm_fp4_sm103.py` - `ruff format --check tests/gemm/test_mm_fp4_sm103.py` - `git diff --check` ## 🧪 Tests - [x] Tests have been updated as needed. - [x] All targeted tests are passing. ```text pytest -q tests/gemm/test_mm_fp4_sm103.py -vv 2 passed, 2 warnings in 5.70s ``` Tested on NVIDIA B300 / SM103. ## Reviewer Notes The runner index is intentionally no longer asserted. It is not stored in `profiling_cache`; the test already verifies the runner identity using `ProfilingCacheKey.runner_class_name`. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Tests** * Updated SM103 FP4 GEMM correctness tests to reflect the current profiling cache format. * Preserved validation that the selected runner is `CutlassFp4GemmRunner`. * Removed obsolete checks for a runner ID that is no longer part of the cache entries. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
## Description Bump version to 0.6.17 for release. **Cut point:** `main` at `369d1ac3` (39 commits since `v0.6.16rc5`). > [!NOTE] > `v0.6.16` is not tagged yet — the release is in progress, currently at `v0.6.16rc5`. > The API diff below is therefore taken against `v0.6.16rc5` (the current contents of > `release-v0.6.16`) rather than the last stable tag, so it shows only the surface that > is new in 0.6.17. The `v0.6.15.post1` → `0.6.16` surface was already reviewed in flashinfer-ai#4142. ## Related Issues (Gated-by PRs) https://github.com/flashinfer-ai/flashinfer/issues?q=is%3Aopen+label%3Av0.6.17 ## Reviewer Notes **API changes review** API changes since v0.6.16rc5, using `scripts/list_apis.sh` ```diff diff -u \ <(scripts/list_apis.sh -d -p --ref v0.6.16rc5) \ <(scripts/list_apis.sh -d -p) --- /dev/fd/11 2026-07-30 16:28:01 +++ /dev/fd/12 2026-07-30 16:28:01 @@ -291,6 +291,13 @@ cp_size: int, enable_pdl: Optional[bool] = None, ) -> tuple[torch.Tensor, torch.Tensor]: +class MixedCommHandler: + @flashinfer_api + def checkpoint_prepare(self): + + @flashinfer_api + def checkpoint_restore(self, comm_backend: CommBackend): + [Global Functions] @flashinfer_api def run_mixed_comm( @@ -458,7 +465,51 @@ hidden_size: int, dtype: torch.dtype, ) -> torch.Tensor: +class UlyssesCommunicator: + @flashinfer_api + def __init__( + self, + group: Optional[ProcessGroup] = None, + *, + max_elems: int, + dtype: torch.dtype, + backend: str = "auto", + device: Optional[Union[torch.device, str, int]] = None, + ): + + @flashinfer_api + def scatter_heads(self, x: torch.Tensor) -> torch.Tensor: + + @flashinfer_api + def gather_heads(self, x: torch.Tensor) -> torch.Tensor: + [Global Functions] +@flashinfer_api +def init_ulysses_a2a( + out_ipc_ptrs: List[int], + signal_ipc_ptrs: List[int], + rank: int, + world_size: int, + full_nvlink: bool, +) -> int: + + +@flashinfer_api +def dispose_ulysses_a2a(fa: int) -> None: + + +@flashinfer_api +def ulysses_a2a( + fa: int, + inp: torch.Tensor, + out: torch.Tensor, + B: int, + S_local: int, + H: int, + D: int, + mode: int, +) -> None: +[Global Functions] @flashinfer_api(trace=concat_mla_k_trace) def concat_mla_k( k: torch.Tensor, @@ -1209,6 +1260,35 @@ ) -> Union[List[torch.Tensor], torch.Tensor]: +@flashinfer_api(trace=trtllm_fp8_per_tensor_scale_routed_moe_trace) +def trtllm_fp8_per_tensor_scale_routed_moe( + topk_ids: torch.Tensor, + routing_bias: Optional[torch.Tensor], + hidden_states: torch.Tensor, + gemm1_weights: torch.Tensor, + output1_scales_scalar: torch.Tensor, + output1_scales_gate_scalar: torch.Tensor, + gemm2_weights: torch.Tensor, + output2_scales_scalar: torch.Tensor, + num_experts: int, + top_k: int, + n_group: Optional[int], + topk_group: Optional[int], + intermediate_size: int, + local_expert_offset: int, + local_num_experts: int, + routed_scaling_factor: Optional[float], + use_routing_scales_on_input: bool, + routing_method_type: int = 0, + do_finalize: bool = True, + enable_pdl: Optional[bool] = None, + tune_max_num_tokens: int = 8192, + activation_type: int = ActivationType.Swiglu.value, + routing_replay_out: Optional[torch.Tensor] = None, + output: Optional[torch.Tensor] = None, +) -> Union[List[torch.Tensor], torch.Tensor]: + + @flashinfer_api(trace=trtllm_fp8_block_scale_moe_trace_dispatch) def trtllm_fp8_block_scale_moe( routing_logits: torch.Tensor, @@ -1686,6 +1766,7 @@ w1_bf16: torch.Tensor, w2_bf16: torch.Tensor, *, + variant=None, num_local_experts: int, hidden_size: int, intermediate_size: int, @@ -1694,6 +1775,12 @@ ) -> Dict[str, torch.Tensor]: + hidden_states_bf16: torch.Tensor, + *, + variant, +) -> Tuple[torch.Tensor, Optional[torch.Tensor]]: + + x: torch.Tensor, ) -> Tuple[torch.Tensor, torch.Tensor]: @@ -2527,6 +2614,7 @@ backend: Literal["cute"] = "cute", out: Optional[torch.Tensor] = None, out_dtype: Optional[torch.dtype] = None, + is_gated: bool = False, ) -> torch.Tensor: [Global Functions] @flashinfer_api @@ -2541,6 +2629,7 @@ backend: Literal["cute"] = "cute", out: Optional[torch.Tensor] = None, out_dtype: Optional[torch.dtype] = None, + is_gated: bool = False, ) -> torch.Tensor: [Global Functions] @flashinfer_api(trace=recurrent_kda_trace) @@ -2571,11 +2660,10 @@ @flashinfer_api def checkpointing_ssu( state: torch.Tensor, - old_x: torch.Tensor, - old_B: torch.Tensor, - old_dt: torch.Tensor, - old_cumAdt: torch.Tensor, - cache_buf_idx: torch.Tensor, + x_cache: torch.Tensor, + B_cache: torch.Tensor, + dt_cache: torch.Tensor, + ring_start: torch.Tensor, prev_num_accepted_tokens: torch.Tensor, x: torch.Tensor, dt: torch.Tensor, @@ -2596,6 +2684,11 @@ cu_seqlens: Optional[torch.Tensor] = None, max_seqlen: Optional[int] = None, enable_pdl: bool = False, + cb_scaled: Optional[torch.Tensor] = None, + cumAdt_vec: Optional[torch.Tensor] = None, + cb_old: Optional[torch.Tensor] = None, + precompute_heads_per_cta: int = 0, + algorithm: str = "auto", ) -> torch.Tensor: [Global Functions] @flashinfer_api(trace=selective_state_update_trace) @@ -2870,11 +2963,13 @@ cute_dsl_impl: str, sinks: Optional[Union[List[torch.Tensor], Tuple[torch.Tensor, ...], torch.Tensor]], + enable_dcp: bool = False, ): cute_dsl_impl: str, sinks: Optional[Union[List[torch.Tensor], Tuple[torch.Tensor, ...], torch.Tensor]], + enable_dcp: bool = False, ): @@ -2898,6 +2993,7 @@ candidate_max: int, cute_dsl_impl: str, sinks: Optional[Union[List[torch.Tensor], Tuple[torch.Tensor, ...], torch.Tensor]], + enable_dcp: bool = False, ) -> int: @@ -2910,10 +3006,25 @@ device: torch.device, cute_dsl_impl: str, sinks: Optional[Union[List[torch.Tensor], Tuple[torch.Tensor, ...], torch.Tensor]], + enable_dcp: bool = False, ) -> Tuple[int, ...]: + *, query: torch.Tensor, + backend: str, + sinks: Optional[List[torch.Tensor]], + cum_seq_lens_q: Optional[torch.Tensor], + max_q_len: Optional[int], + return_lse: bool, + enable_dcp: bool, + cp_world: int, + cp_rank: int, + causal_seqlens_kv_global: Optional[torch.Tensor], +) -> str: + + + query: torch.Tensor, out_dtype: torch.dtype, bmm1_scale: Union[float, torch.Tensor], bmm2_scale: Union[float, torch.Tensor], @@ -2928,10 +3039,17 @@ cute_dsl_impl: str = "auto", cum_seq_lens_q: Optional[torch.Tensor] = None, max_q_len: Optional[int] = None, + enable_dcp: bool = False, + cp_world: int = 1, ) -> Optional[str]: buckets: tuple[int, ...], num_pages: int, profile_seq_len: int, + has_sparse_mla_top_k_lens: bool = False, + sparse_top_k_width: int = 0, + enable_dcp: bool = False, + cp_world: int = 1, + cp_rank: int = 0, ) -> TuningConfig: [Global Functions] @@ -2963,6 +3081,11 @@ cum_seq_lens_q: Optional[torch.Tensor] = None, max_q_len: Optional[int] = None, multi_ctas_kv_counter_buffer: Optional[torch.Tensor] = None, + sparse_mla_top_k_lens: Optional[torch.Tensor] = None, + enable_dcp: bool = False, + cp_world: int = 1, + cp_rank: int = 0, + causal_seqlens_kv_global: Optional[torch.Tensor] = None, ) -> Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]]: @@ -4249,7 +4372,7 @@ @flashinfer_api def nvfp4_quantize_cute_dsl( input: torch.Tensor, - global_scale: torch.Tensor, + global_scale: float | torch.Tensor, sf_layout: int = SF_LAYOUT_128x4, enable_pdl: bool | None = None, ) -> Tuple[torch.Tensor, torch.Tensor]: @@ -4752,6 +4875,7 @@ tie_break: int = TopKTieBreak.NONE, dsa_graph_safe: bool = False, row_starts: Optional[torch.Tensor] = None, + page_table_row_starts: Optional[torch.Tensor] = None, ) -> torch.Tensor: ``` <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Chores** * Updated the application version from 0.6.16 to 0.6.17. <!-- end of auto-generated comment: release notes by coderabbit.ai --> Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
Re-attaches the ROCm halves that #322 evacuated into benchmarks/routines/rocm/, against a routines/attention.py upstream changed by +1991/-177. Three of the port's four seams shrank because upstream converged on them: - flashinfer_benchmark.py no longer needs load_routine_group. Upstream now imports each routine module lazily inside run_test(), which is what the shim existed to do. The helper stays (tests/rocm/test_benchmark_harness.py covers it) but the driver no longer calls it, so the fork's delta on that file is now just the csv.writer and add_timing_budget_args. - routines/attention.py: "auto" is already in the --backends choices and already admitted by the decode and paged-prefill wrapper branches. What the port still has to say is that on ROCm `auto` needs NHD, no graph capture and no tensor cores, or the wrapper silently resolves to fa2 and the row measures fa2 while labelled auto. - The MLA routine now records backend resolution too; it did not before, and MLA is one of the paths that defaults to AITER. bench_timing_kwargs keeps passing the l2_flush* trio even though upstream deprecated it in favour of cold_l2_cache. Not an oversight: the cold_l2_cache path sizes the flush buffer at 2x props.L2_cache_size, which on CDNA reports the 4 MB L2 rather than the 256 MB Infinity Cache, so it would leave the LLC resident and measure warm. Unverified on-box at this commit -- confirmed in the container run before merge. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Two defects the merge introduced, both found by running rather than reading.
flashinfer/quantization/ is a package in v0.6.18, and its __init__ imports
fp8_quantization and fp4_quantization eagerly. Those reach for CUDA-only jit
exports at import time, so `import flashinfer` died on ROCm:
flashinfer/rocm/prefill.py:39 from ..quantization import packbits
-> flashinfer/quantization/__init__.py:29
-> flashinfer/quantization/fp4_quantization.py:34
ImportError: cannot import name 'sm121a_nvcc_flags' from 'flashinfer.jit'
Importing the submodule directly does not help -- the package __init__ runs
either way -- so the gate goes where the coupling is, matching comm/__init__.py.
packbits is the only member ROCm compiles. __all__ is split the same way, so a
star-import on ROCm advertises only what is bound.
scripts/git_describe_rocm.py picked the closest ancestor tag regardless of
shape. After the merge that is upstream's own v0.6.18:
v0.6.18 ancestor=Y distance=364
v0.5.3+amd.2 ancestor=Y distance=1249
$ python3 scripts/git_describe_rocm.py
v0.6.18-364-g3b742fa08
which setuptools-scm turns into a version with no +amd segment -- an
amd-flashinfer wheel indistinguishable from upstream's, with nothing to say so.
Fork tags now win at any distance; plain tags remain the fallback for a tree
that has none, which is what the existing 13 cases exercise. Restricting the
`git tag -l` glob to `v*+amd.*` instead would have broken six of them.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
include/ goes from 214 to 287 upstream headers. The wheel filter is unchanged (_WHEEL_HEADER_FILES still names flashinfer/fp16.h, byte-identical across the merge base, HEAD and v0.6.18), but MANIFEST.in grafts include/ wholesale, so the sdist grows by the full upstream delta while the wheel does not. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Catch-up with 128f9df (#343). Auto-merged; the only overlap is pyproject.toml, where #343 rewrote the `slow` marker description and this branch dropped a testpaths entry for a test v0.6.18 deletes. A rebase is not usable here: the branch carries the v0.6.18 merge commit, so replaying it would re-resolve all 24 conflicts against the same upstream tag for no gain. The v0.5.3 upgrade took a catch-up merge for the same reason (61c11fd). Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…ream
`pre-commit run -a` on the merge rewrote 24 upstream files and left 7 mypy
errors. Both are merge fallout, and neither is fixed by committing the
rewrites -- every reformatted upstream file is a permanent merge conflict.
pyproject.toml gains upstream's [tool.ruff] extend-exclude verbatim. The port
kept its own pyproject through the merge, so it did not inherit the block
v0.6.18 added to stop ruff-format rewriting generated and vendored kernels;
without it, 16 upstream files reformat.
markdownlint is scoped to the eight markdown files the port owns plus the four
it overrides. It is a fork-added hook -- upstream does not run it -- so aiming
it at the whole tree rewrote docs/design_docs/, .claude/skills/add-cuda-kernel
and flashinfer/moe_ep/**/*.md.
mypy, four sites:
- flashinfer/quantization is a package now, so `from ..quantization import
packbits` is ambiguous between the submodule and the re-exported function.
Runtime picks the function; the two ROCm callers now say so explicitly.
- JitSpec became an ABC, so is_aot/write_ninja moved to JitSpecNvcc.
refresh_aiter_jitspec keeps the JitSpec annotation gen_jit_spec declares and
carries two targeted ignores, rather than widening the delta on jit/core.py.
- The IS_HIP arm rebinds __version__, which upstream's IS_CUDA arm now binds
from three places.
Six files also take genuine formatting fixes from the hooks -- all lines the
re-indentation under `if IS_CUDA:` pushed past the line limit.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
… tests
Findings from /code-review on the restore, each reproduced on gfx942 before
being fixed.
benchmarks/routines/attention.py gained three module-level imports in v0.6.18
that abort the whole harness on ROCm, before any benchmark runs:
ImportError: cannot import name 'autotune' from 'flashinfer'
`autotune` is bound only in the IS_CUDA arm; `flashinfer.fp4_quantization`
pulls CUDA-only jit exports; `flashinfer.prefill` is shadowed to
flashinfer.rocm.prefill, which has no trtllm_fmha_v2_prefill. Every use is on a
CUDA-only path -- NVFP4 KV, the trtllm-fmha-v2 backend, and --autotune /
--autotune_cache -- so all three move to their call sites. This had made every
ROCm hook restored in 3b742fa unreachable.
test_benchmark_harness.py::test_hip_gqa_group_sizes_match_the_kernel_dispatch
read include/flashinfer/utils.cuh -- upstream's header, which v0.6.18 gave a
`group_size == 6` arm. The HIP decode kernel compiles against the forked
rocm/dispatch.cuh, still {1,2,3,4,8}, so HIP_DECODE_GQA_GROUP_SIZES was right
and the test was reading the wrong file. It has pointed at the upstream
original since the headers were forked; the merge is only what exposed it.
test_jit_env.py::test_nvshmem_helpers_stay_absent_on_rocm named
gen_nvshmem_module(), which v0.6.18 deletes -- the NVSHMEM build folded into
gen_mixed_comm_module(). That renders .cu templates into FLASHINFER_GEN_SRC_DIR
*before* it reaches `import nvidia.nvshmem`, so the "builds nothing on the way"
property the test guarded no longer held either. flashinfer.comm.mixed_comm
joins CUDA_ONLY_MODULES and the test asserts the gate instead.
flashinfer_benchmark.py gets load_routine_group back. Upstream's lazy imports
cover the import-deferral, but not the diagnostic: without the loader a
CUDA-only routine group surfaces a raw ImportError rather than being named
unavailable, and test_runner_routes_each_group_through_the_loader pins the
wiring.
Declined: extending CUDA_ONLY_MODULES to the other 12 comm submodules v0.6.18
added. dcp_alltoall, allreduce and trtllm_moe_alltoall already raise the gate's
ImportError transitively, and quantized_allreduce and ulysses import cleanly on
ROCm in this image -- gating them without a reproduced failure is speculative.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
There was a problem hiding this comment.
Pull request overview
Syncs the AMD ROCm port of FlashInfer to upstream v0.6.18, bringing the repo’s Python/C++ surface area and test suite into alignment with the version pinned by vLLM/SGLang, while re-applying ROCm-specific routing, packaging, and CI adjustments needed to keep the port functional.
Changes:
- Updates/extends multiple CUDA and ROCm-facing modules (JIT, comm, attention/norm/sampling/topk, versioning) to match upstream v0.6.18 structure and APIs.
- Adds/updates test coverage across trace/reference correctness, sharding infrastructure, ROCm utilities, and various CUDA-only paths.
- Refreshes CI/devcontainer and repository automation scripts/config to match the new upstream layout and workflows.
Reviewed changes
Copilot reviewed 65 out of 2221 changed files in this pull request and generated 2 comments.
Show a summary per file
| File | Description |
|---|---|
flashinfer/version.py |
Updates version metadata loading and backward-compat aliasing for git commit/version fields. |
tests/utils/test_jit_warmup.py |
Adds a CUDA-arch-specific warmup test for SM90-only fa3 backend. |
csrc/bgmv_moe/moe_bgmv_ops.h |
Introduces a public header for BGMV MoE TVM-FFI entry points. |
tests/comm/test_nvshmem.py |
Adjusts NVSHMEM smoke checks and import paths for basic API access. |
tests/rocm/test_git_describe.py |
Adds regression coverage for fork-tag precedence over nearer upstream tags. |
tests/rocm/test_benchmark_harness.py |
Fixes ROCm header reference for HIP group-size dispatch assertions. |
flashinfer/rocm/api.py |
Moves ROCm packbits exports to quantization.packbits submodule to avoid eager CUDA-only imports. |
flashinfer/rocm/prefill.py |
Updates ROCm imports to avoid quantization package side effects and keep ROCm importable. |
.github/workflows/pre-commit.yml |
Switches workflow Python selection to .python-version and adds CUDA config validation step. |
ci/docker-tags.yml |
Updates CI image tags to the newer CUDA CI baselines. |
docs/api/*.rst |
Adds/expands API docs for new/reshaped modules (topk, pod, mhc, mamba, kda, etc.). |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Upstream widened 10 of the 11 sampling entry points and flashinfer/sampling.py
is shared verbatim, so every ROCm call raised at dispatch:
RuntimeError: sampling::sampling_from_probs() expected at most 6 argument(s)
but received 9 argument(s).
Three separate widenings, none of which the ROCm kernels can use:
* the scalar philox pair became optional per-request seed/offset tensors
(7 ops, +2 each). One scalar seed covers the whole batch here, so an
actual tensor is rejected rather than silently ignored -- the Python
default path yields ints, so only an explicitly tensor-seeded call trips it.
* the five *_from_probs ops gained a `valid` output (+1). ROCm's kernels
carry the last-valid-index fallback but not upstream's reject flag, so
every row samples and `valid` is filled true.
* the renorm ops gained multi-CTA scratch (row_states_buffer) and an
is_deterministic switch for AIR top-p (+1/+2). ROCm keeps the single-CTA
ternary-search kernels, which are deterministic and need no scratch.
top_k_top_p_sampling_from_{logits,probs} also grew a radix top-k fast path that
builds csrc/topk.cu; csrc/rocm has no such file, so it failed in ninja rather
than at import. Gated off on HIP.
Measured on gfx950 (MI350X, ROCm 7.15, torch 2.12): tests/rocm/test_sampling.py,
test_logits_processor.py and test_quantization_segments.py go from 6 collected
failures plus a hard dispatch error on the first sampling call, to fully green.
test_sampling_binding_abi.py reads both sides statically -- it needs no GPU and
would have caught all three widenings at merge time. test_quantization_segments
follows _fake_packbits into the packbits submodule, where 0.6.18 left it when
quantization became a package.
Matching the new ABI on arity was not enough. v0.6.18 replaced the
`probs = probs.float()` in the top_k_renorm_probs and top_k_mask_logits
wrappers with an assert that *permits* fp16 and bf16 and passes the tensor
through untouched, allocating the output with torch.empty_like. ROCm's
kernels are float32-only and the binding still did static_cast<float*> with
vocab_size as the row stride, so a half input was read -- and the overrun
written back -- across twice the bytes the allocation holds. No crash, no
wrong-dtype error: silent corruption of both buffers.
It is reachable from a shipped path, not just a hand-written call.
LogitsTopKOp passes tensor.data straight through, so LogitsPipe([TopK(),
Softmax(), Sample()]) on bf16 logits -- the natural dtype off a model head --
hits it, as does top_k_top_p_sampling_from_logits with the default
filter_apply_order. The previous commit made this worse: before it, the same
call died on the arity mismatch.
A sweep of all 11 wrappers confirms these two are the only ones that lost the
cast; every other op still forces float32 in Python.
Two more from the same review:
* top_p_renorm_probs sizes the AIR top-p scratch from batch and vocab and
allocates it fresh per call -- ~12 MB at batch 256 / vocab 152k. ROCm
never reads it, and TopPOp routes through this API, so a decode loop paid
it every step.
* tests/attention/test_page.py stopped collecting entirely: 0.6.18 gave it a
top-level nvfp4 import, so a bare `pytest` -- what CI runs -- died during
collection, which a tests/rocm/-only run cannot see. Dropped from
testpaths; its one HIP-relevant case, test_append_paged_kv_cache, is
covered further by tests/rocm/test_append_paged_kv_cache_aiter.py. The
four quantization module names join CUDA_ONLY_MODULES so a direct import
reports "CUDA-only" rather than a bare sm121a_nvcc_flags ImportError.
Verified on gfx950: the three new sampling cases and the import-gate,
dispatch, quantization and binding-ABI tests are green (58 cases), and
`pytest --collect-only` is now clean across every remaining testpath. The
unguarded version was not run -- it corrupts memory, and both nodes are
carrying live suites.
Three findings from reviewing d023ac8, two of which it got wrong. The fp32 guard was the wrong shape. At 128f9df the wrappers upcast fp16/bf16 themselves, so LogitsPipe([TopK()]) on bf16 model logits -- the ordinary case -- worked on ROCm; a TORCH_CHECK turned that into a hard failure on ROCm while it kept working on CUDA, and made sampling.py's "Supported dtypes: float32, float16, bfloat16" false here. The binding now upcasts and casts the result back, which keeps the overrun protection without the capability loss. The test follows: it asserts the dtype round-trips and matches the fp32 result exactly (half -> fp32 is lossless), rather than asserting a raise. A/B: dropping the copy_ back leaves the caller's uninitialized buffer and fails all four cases. The four new CUDA_ONLY_MODULES entries never took effect. gate_cuda_only_modules had one call site, flashinfer/comm/__init__.py, and `import flashinfer` does not reach flashinfer.comm on HIP -- so the gate was installed only for callers that imported comm first, and python -c "import flashinfer; import flashinfer.quantization.fp4_quantization" still raised the bare sm121a_nvcc_flags ImportError the commit claimed to replace. tests/rocm/test_comm_import_gate.py could not catch it: its first test imports flashinfer.comm, installing the gate before the parametrized ones run. Called from flashinfer/rocm/api.py now, ahead of its own imports. Verified: both names report "is CUDA-only and not available on ROCm". Also restores one shape the testpath deletion left uncovered. Every 5-D combined cache test is @requires_aiter, so on a box without the package nothing exercised the native kernel on that layout -- which is what the deleted test_append_paged_kv_cache did. Added the native-only case beside it. The AIR workspace shortcut moved above the sizing arithmetic it was skipping, so the dead ~12 MB computation no longer runs at all. gfx950: 8 targeted cases plus the sampling, logits-processor, quantization, import-gate, dispatch and binding-ABI suites green; pytest --collect-only clean across every testpath.
Third review round on the sampling work; all findings taken. top_p_renorm_probs was the one op left with no dtype handling in C++, safe only because its Python wrapper still calls probs.float() -- the very line v0.6.18 deleted from the two siblings, which is what caused the corruption this series is fixing. It goes through as_fp32 as well now, so the next sync that drops the cast there cannot reopen it silently. as_fp32 also validates before it branches. The fp32 fast path handed the caller's buffer straight to the kernel with no CHECK_INPUT and no shape check, while the cast path tolerated any layout via new_empty + copy_ -- so a non-contiguous fp32 out scribbled past the end where the identical fp16 out was accepted. Unreachable from Python (torch.empty_like), but the asymmetry was new in a8523a3. Nothing in the suite failed if the gate_cuda_only_modules() call was deleted again: every existing gate test either calls it directly or depends on an earlier test in the same process having imported flashinfer.comm, which under -n auto need not even be the same xdist worker. The new case runs a subprocess that imports flashinfer alone and asserts flashinfer.comm is absent from sys.modules before checking. A/B: commenting out the call reproduces the sm121a_nvcc_flags ImportError and fails the test. The top_p workspace shortcut goes back to the two-line `if IS_HIP: ws_size = 1`. The early return read better but duplicated an upstream dispatch call whose argument list just changed in this very sync, and it saved nothing -- the old form already allocated one byte. Also: two docstrings claimed more than they proved (the dtype round-trip does not prove the cast-back ran, torch.equal does; upstream's test_page.py did cover the 5-D layout, it just asserted nothing about the bytes), a scatter comparison used default bf16 tolerance instead of rtol=0, and a comment in test_rocm_dispatch.py still said the gate might not be up yet.
The full gfx942 suite at fe3c4b4 came back with 11 failures, in two files, both host-only and neither GPU-dependent. Nine are the ninja generator tests. Their stub jit_env is a SimpleNamespace carrying only INCLUDE_DIR and CSRC_DIR, and generate_ninja_build_for_op started reading FLASHINFER_JIT_DIR when it moved to absolute output paths for v0.6.18 (365533b) -- so every test that renders a ninja file raised AttributeError. A fixture gap introduced by that commit, not a product defect. Two are the coverage manifest: it redirects flashinfer/quantization.py, which v0.6.18 replaced with a package. Retargeted at flashinfer/quantization/packbits.py, the member that actually wraps csrc/rocm/quantization.cu -- the package __init__ carries a fork diff now and needs no redirect. A sweep of every redirect_owned and unowned entry against the tree finds no other stale path. Both files green on gfx950 (125 cases).
Fourth review round; all findings were about tests that passed for the wrong reason rather than about the code under them. test_every_row_reports_valid asserted valid.all() on a buffer the wrapper allocates with torch.empty, so it held whenever the caching allocator handed back a block with nonzero bytes -- which is the normal case once anything else in the file has run in the same worker. It only discriminated on a cold process, i.e. a coin flip in the passing direction under -n auto. It now calls the raw op with a false-filled valid. test_a_tensor_seed_is_rejected covered 1 of the 7 reject_per_request_seed call sites. A forgotten site on the next sync collapses every row onto one philox stream, which is invisible in the output -- the samples are still valid tokens. Parametrized over the six public entry points. test_gate_tolerates_a_foreign_marker was the last test reaching the finder's insert branch, and installing the gate in rocm/api.py made it fall through to the real finder and widen that instead. So changing insert to append -- which puts _CudaOnlyFinder behind PathFinder and silently stops the gate working in a fresh process -- would have shipped green. It lifts the real finder out for the duration now. Also: as_fp32 upcast any dtype, including float64 and integers, so the C++ guard depended on the Python assert it was written to outlive; and top_p's ignored workspace is pinned at numel()==1 so a future kernel that reads it fails loudly rather than overrunning one byte. Two manifest reasons that v0.6.18 falsified are corrected.
…nned The guard added in 64bf283 did not work. A/B'd it by changing the finder's sys.meta_path.insert(0, ...) to append -- which puts _CudaOnlyFinder behind PathFinder and silently stops the gate working -- and every gate test stayed green, including the one whose docstring claimed to catch exactly that. The reason is that all of them probe names like flashinfer.comm._not_a_real_module, which PathFinder cannot resolve either, so the finder's position never matters. The new test gates flashinfer.quantization.packbits, pops it from sys.modules so the cached object cannot short-circuit the lookup, and restores it afterwards. A/B: insert -> append now fails it with DID NOT RAISE. Lifting the real finder out in test_gate_tolerates_a_foreign_marker is kept -- it does restore that test's reach into the install branch -- but its docstring no longer claims to cover the ordering.
v0.6.18 brings csrc/fmha_v2/ and flashinfer/jit/attention/fmha_v2/, 61 files carrying the NVIDIA TensorRT Source Code License, which forbids redistribution. Neither path existed at the merge base, so this PR is what introduces them: $ git ls-tree -r --name-only origin/amd-integration -- csrc/fmha_v2 | wc -l 0 Measured on the built artifacts rather than reasoned about. The sdist carried 77 csrc/fmha_v2 entries, and the wheel carried four Python modules including the licensed generator_utils.py -- Copilot's review flagged the sdist case and understated it. The wheel needed exclude-package-data, not the packages.find exclude I reached for first. That exclude does work (setuptools resolves the package list without fmha_v2), but the directory has no __init__.py, so the files arrive as package *data*: include-package-data defaults true for a pyproject project, and setuptools-scm's file finder then sweeps every git-tracked file under flashinfer/ into the wheel. After: wheel 986 -> 982 entries, sdist 2101 -> 2027, and a scan of both for the licence string returns nothing but this commit's own comments naming it. The installed wheel still imports, with the module absent. The three loose csrc/*fmha_v2*.cu files stay -- they are Apache-2.0, only the directory is not. Separately, the pre-commit workflow's new "Validate CUDA configuration" step is dropped. It is not merely noisy here, it is unsatisfiable: with pyproject.toml reshaped to suit its hand-rolled TOML parser it then demands ERROR: pyproject.toml cu13 extra must contain 'nvidia-cutlass-dsl[cu13]>=4.6.2a0' which a ROCm wheel cannot declare. It is what turned the pre-commit lane red.
…issed Review of 7bded9b found its central claim backwards, and I had measured it backwards too. exclude-package-data does nothing here; the packages.find exclude alone is what drops the files. A/B on this tree, that config only: packages.find exclude ONLY -> entries 982 fmha_v2 0 My earlier sequential reading -- 4 files with the packages.find exclude, 0 after adding exclude-package-data -- was confounded by a stale build/lib. The mechanism is the reverse of what the comment said: `find` defaults to namespaces = true, so flashinfer/jit/attention/fmha_v2 is discovered as a namespace package despite having no __init__.py, and its files ship as package *modules*, which exclude_package_data never sees. The inert block is removed rather than left as a second net that does not catch anything. csrc/cudnn_sdpa_utils.h is the same licence class -- LicenseRef-NvidiaProprietary, "disclosure or distribution ... strictly prohibited" -- and shipped in the sdist. It predates v0.6.18, which is why the two new prunes missed it; a scan for both markers rather than just the TensorRT one finds it and nothing else. release.yml runs the same CUDA validator on a `pull_request:` trigger with no branches filter, so it fires on PRs to amd-integration whenever one of its watched paths changes -- and fails for the reason the pre-commit step was dropped. Removed there too. The script stays: pr-test.yml, release-ci-docker.yml and nightly-release.yml are gated on main/cron and still use it. The pre-commit comment named the CUDA-extras requirement, but the line-based TOML reader fails first on this file's multi-line `dev = [`, so someone acting on it would add the extras and still get exit 1. Corrected. test_redistribution_licences.py is the durable guard. The existing artifact test cannot be one: test_build_backend.py builds in a non-git tree, so setuptools-scm finds no files and the wheel omits these paths whether or not the exclusions exist. Not fixed, and deliberate: flashinfer/jit/attention/modules.py still ships and imports .fmha_v2 at module level, so it raises ModuleNotFoundError on a CUDA host. This is a ROCm-only wheel whose CUDA arms are dead by fork policy, and excluding modules.py would drop code the fork does not own for no gain here.
The file spells out 'TensorRT Source Code License' and 'LicenseRef-NvidiaProprietary' in _MARKERS, and it is a tracked .py, so the scan flagged the guard itself the moment it was committed. Caught by the A/B restore, not by the first run -- before the commit the file was untracked and git ls-files did not list it.
Review of 8a9da5f found the guard I had just written passing for the wrong reasons in four ways. All measured against probe files, not argued. The hand-kept _EXCLUDED_PREFIXES was the root of two of them and is gone: the covered set is now parsed out of MANIFEST.in's own `prune`/`exclude` lines. Before, appending a prefix to the tuple turned the test green while nothing was excluded from either artifact -- the one-line "fix" a future sync would reach for. Splitting `prune` (directory, startswith) from `exclude` (exact path) also fixes csrc/cudnn_sdpa_utils.h being consumed by startswith, which silently covered any sibling sharing the prefix. The suffix allowlist missed .jinja (65 tracked files), .inl and .cc, all of which ship in the sdist. Dropped entirely -- read_text(errors="ignore") already tolerates binaries, and an allowlist is precisely the thing that stops covering a new file type without saying so. A/B, with both holes probed at once: AssertionError: these carry a no-redistribution licence and no MANIFEST.in rule covers them: ['csrc/demandal_ab_probe.hpp', 'csrc/demandal_ab_probe.jinja'] The workflow check named two files and matched a literal string. It now globs and parses the trigger, so a workflow counts only if a push/pull_request event has no `branches:` filter -- which is the rule the PR already follows for the dormant upstream workflows, rather than a list to keep in step. That surfaced pr-test, nightly-release and release-ci-docker as carrying the same step; they stay untouched because they are gated on main or on a schedule against the default branch, and editing a dormant upstream file buys a permanent conflict. The wheel side gets its own case: MANIFEST.in governs the sdist only, so a restricted file under flashinfer/ must also match a packages.find exclude.
There was a problem hiding this comment.
🔵 Needs a closer look
Some newly-modified tests can raise at collection time on CPU-only/no-GPU environments due to unguarded torch.cuda.get_device_capability() calls.
Review details
- Files reviewed: 66/2234 changed files
- Comments generated: 0 new
- Review effort level: Lite
demandal25
added a commit
that referenced
this pull request
Sep 2, 2026
Record the v0.6.18 merge that #344's squash discarded
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
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.
Summary
Syncs the port from upstream
v0.5.3tov0.6.18— 1125 upstream commits and 2224 files since the merge base2628bebcf(2025-11-20). vLLM and SGLang both pinflashinfer-python==0.6.18exactly on main, so this is the surface their call sites are written against.Scope is deliberately the sync itself: land the merge, restore the ROCm halves, get the port green. No new upstream features are adopted, the forked headers are not refreshed, ROCm's public API is not brought to parity with 0.6.18's reworked wrappers, and the AITER pin stays at 0.1.20. The two follow-ups that scoping implies are written down at the bottom rather than left implicit.
Read it commit by commit:
ad239a313is the merge with every shared code file taken from upstream verbatim, and each commit after it restores one subsystem.What changed
The merge —
ad239a313git merge-treeraised 24 conflicts, 15 of them code — modest for a jump this size, and the payoff of the evacuation work in #305, #314, #318, #325 and #334. Resolution was upstream-verbatim for every shared code file, sogit diff v0.6.18over those paths is empty at that commit and the ROCm side comes back in reviewable pieces afterwards.flashinfer/norm.py→flashinfer/norm/— upstream turned the module into a package, so the deletion is accepted and the dispatch is re-homed infc730dd08. This is the one conflict that fails silently if fumbled: leavingnorm.pybeside the package resolves to the package and the ROCm routing just disappears, with no import error.CHANGELOG.mdandtests/utils/test_block_sparse_indices_to_vector_sparse_offsets.py(both gone upstream);version.txtstays deleted, the port versions from setuptools-scm.README.md,CONTRIBUTING.md,benchmarks/README.md,CLAUDE.md,.claude/skills/benchmark-kernel/SKILL.md,pyproject.toml..gitignoreis the union.pytest.iniis not a conflict — it predates the merge base and upstream has not touched it since, so the port's deletion carries through and[tool.pytest.ini_options]still governs test selection.Restoring the ROCm side
flashinfer/jit/{env,core}.py(365533b3e) — the two files that gate everything else.env.pyis unchanged in shape; two things shrank, because upstream deletedget_nvshmem_{include,lib}_dirsoutright and now does theextra_include_paths→Pathconversion itself.core.pyneeded real work: upstream refactoredJitSpecinto an ABC with aJitSpecNvccsubclass and movedbuild_and_load()onto the base, so the port's oldclass_name=plumbing no longer fits — it is dropped rather than ported, since no caller has ever passed it.flashinfer/jit/rocm/cpp_ext.py(same commit) — emit absolute output paths. Upstream moved its own generator to absolute paths and switchedbuild()to run ninja with-C self.build_dirinstead of the JIT root; the ROCm generator still emitted$name/<obj>, which under the new workdir would have nested every artifact one level deeper.__init__files (275f64be8) — mechanical. Upstream's export list re-indented underif IS_CUDA:, the port'selif IS_HIP:/else:arms re-attached unchanged. No symbol the port binds was renamed or dropped.flashinfer/{activation,page,pod,rope}.py(2d1228dca) — thebackend="auto"|"native"|"aiter"parameter and itsmaybe_*short-circuit.pod.pyabsorbs a fifth CUDA-onlyplan()argument (uniform_q_len) as one more element of the tail the port already splits off;rope.pygrew a seventh_fake_*op.flashinfer/norm/__init__.py(fc730dd08) — plus one addition, below.3b742fa08,4b20e8011) — against aroutines/attention.pyupstream changed by +1991/−177.Defects this surfaced
import flashinferwas dead on ROCm (a5fc63be4).flashinfer/quantization/is a package in 0.6.18 and its__init__imports fp4/fp8 eagerly, which reach CUDA-only jit exports:ImportError: cannot import name 'sm121a_nvcc_flags'. Importing the submodule directly does not help — the package__init__runs either way — so the gate goes where the coupling is, matchingcomm/__init__.py.packbitsis the only member ROCm compiles.a5fc63be4).scripts/git_describe_rocm.pytook the closest ancestor tag regardless of shape; after the merge that is upstream's ownv0.6.18at distance 364 againstv0.5.3+amd.2at 1249, and a bare tag yields a version with no+amdlocal segment and no error to say so. Fork tags now win at any distance. Restricting thegit tag -lglob tov*+amd.*instead would have broken six of the thirteen existing cases.nvidia-cutlass-dslis ever installed (fc730dd08). Upstream routes norm through CuTe DSL by default and ROCm escaped only via anexcept ImportError— i.e. because the package happens to be absent._USE_CUDA_NORMis now forced underIS_HIP, so the default the port chose in fix(rocm): default the RMSNorm ops to the native kernel, with the benchmark that decided it #332 does not depend on what is in site-packages.4b20e8011). Three module-level imports added in 0.6.18 fail on ROCm, soparse_argsdied onImportError: cannot import name 'autotune' from 'flashinfer'— which made every ROCm hook restored in3b742fa08unreachable. All three uses are CUDA-only paths and move to their call sites.fe3c4b4c0).flashinfer/sampling.pyis shared verbatim and 0.6.18 widened 10 of its 11 entry points, so the first call died withsampling::sampling_from_probs() expected at most 6 argument(s) but received 9. Three unrelated widenings landed at once: the scalar philox pair became optional per-request seed/offset tensors (7 ops), the five*_from_probsops gained avalidoutput, and the renorm ops gained multi-CTA scratch plus anis_deterministicswitch for AIR top-p. None maps onto a ROCm kernel — a per-request seed is rejected rather than silently collapsed toseed_val,validis uniformly true because ROCm's kernels carry the last-valid-index fallback but not upstream's reject flag, and the scratch is unused because the single-CTA ternary-search kernels are deterministic already.top_k_top_p_sampling_from_{logits,probs}also grew a radix top-k fast path that buildscsrc/topk.cu, whichcsrc/rocmdoes not have, so it failed inside ninja rather than at import; gated off on HIP.tests/rocm/test_sampling_binding_abi.pycompares the two sides statically — no GPU, and it would have caught all three widenings at merge time.tests/attention/test_page.pystopped collecting, so a barepytestdied before running anything (d023ac824). 0.6.18 gave it a module-level nvfp4 import. It is dropped fromtestpaths— its one HIP-relevant case is covered further bytests/rocm/test_append_paged_kv_cache_aiter.py, and a native-only 5-D case was added there because every existing one is@requires_aiter. Only a full-testpath run sees this; atests/rocm/-only run cannot.a8523a350).gate_cuda_only_modules()had one call site,flashinfer/comm/__init__.py, andimport flashinferdoes not reachflashinfer.common HIP — so the gate was up only for callers that imported comm first. Every existing test hid it by importing comm itself. Called fromflashinfer/rocm/api.pynow.7bded9b66).csrc/fmha_v2/andflashinfer/jit/attention/fmha_v2/are new in 0.6.18 and 61 of their files carry the NVIDIA TensorRT Source Code License, which forbids redistribution. The sdist shipped 77 of them and the wheel four Python modules, including the licensedgenerator_utils.py. The wheel case neededexclude-package-data, not thepackages.findexclude: the directory has no__init__.py, so the files arrive as package data —include-package-datadefaults true for a pyproject project and setuptools-scm's file finder then sweeps every tracked file underflashinfer/in.pre-commitCI was red on the newValidate CUDA configurationstep (7bded9b66). Not noise: withpyproject.tomlreshaped to suit its hand-rolled TOML parser it then demandsnvidia-cutlass-dsl[cu13]>=4.6.2a0in acu13extra, which a ROCm wheel cannot declare. The step is dropped.4b20e8011).test_hip_gqa_group_sizes_match_the_kernel_dispatchread upstream'sinclude/flashinfer/utils.cuh, which gained agroup_size == 6arm; the HIP decode kernel compiles against the forkedrocm/dispatch.cuh, still{1,2,3,4,8}, so the constant was right and the test was reading the wrong header — it has pointed at the upstream original since the headers were forked, and the merge only exposed it.test_nvshmem_helpers_stay_absent_on_rocmnamedgen_nvshmem_module(), which 0.6.18 deletes.What v0.6.18 adds without conflicting
A merge only reports files both sides changed. 1767 files arrive that the port did not have, and several change behaviour with nothing to announce them (
b9cbb433e,07ec7c8f1):pr-labeler.yml(pull_request_target, no branch filter, needsFLASHINFER_BOT_TOKEN) would be red on every PR andissue-claim.yml(cron*/30) roughly 48 times a day. They are dropped.pr-test.yml,pr-api-doc-checks.ymlandpr-api-doc-comment.ymlarebranches: [main]and stay dormant, so they are left alone — deleting an upstream file buys a permanent conflict on every future sync.pre-commit run -awanted to rewrite 24 upstream files. The port kept its ownpyproject.tomlthrough the merge and so never inherited the[tool.ruff] extend-excludeblock 0.6.18 added for generated and vendored kernels; without it, 16 upstream files reformat.markdownlintis fork-added — upstream does not run it — so aiming it at the whole tree rewrotedocs/design_docs/,.claude/skills/add-cuda-kernelandflashinfer/moe_ep/**/*.md. Both fixes are in the config, not the files.flashinfer.comm.mixed_commjoinsCUDA_ONLY_MODULES: it is where the NVSHMEM build landed, and it renders.cutemplates intoFLASHINFER_GEN_SRC_DIRbefore reachingimport nvidia.nvshmem, so without the gate a ROCm caller writes CUDA sources on its way to the failure.Architecture / design notes
Merge, not rebase. The two prior upgrades (#156 for v0.3.1, #173 for v0.5.3) were both
git merge <upstream tag>into a topic branch, and the same reasoning holds: a rebase replays 353 fork commits over 1125 upstream commits and re-resolves the same files at every step, where a merge resolves each once. The catch-up with128f9dfd5mid-branch is a merge for the same reason — rebasing a branch that carries the v0.6.18 merge commit would re-raise all 24 conflicts against the same tag for no gain, which is why the v0.5.3 upgrade did the same thing in61c11fdb8.Both architectures. The merge changes zero files under
include/flashinfer/rocm,csrc/rocm,flashinfer/rocm,tests/rocm,benchmarks/rocmordocs/rocm, and the ROCm jinja templates include onlyflashinfer/rocm/*headers, so no upstream header reaches the HIP build. The 196-fileinclude/replacement is inert for gfx942 and gfx950 alike and the risk is concentrated in Python — which is why the arch-specific treatment below is thin rather than absent.Coverage needs no code change, but does need the tag.
scripts/amd_coverage.pyderives its base from the fork's own*+amd.*tag, so it self-corrects the momentv0.6.18+amd.1is cut. Until then it derivesv0.5.3, which is not an ancestor of the merge, and reports 827 owned.pyfiles instead of 39 — green, because the arch-caps job only assertsownedis non-empty, and meaningless. Cutting the tag right after this merges is the fix; papering over a one-PR window in the tool would be worse. Verified rather than assumed, in a throwaway clone atf58311959: taggingv0.6.18+amd.1moves the version fromv0.5.3+amd.2.dev1266tov0.6.18+amd.1, the derived base to69ff11fc4("bump version to 0.6.18"), and the owned set from 827 files to 56.A user holding a 0.5.3-line
amd-flashinfer-jit-cachewheel will hit aRuntimeErrorfromflashinfer/jit/rocm/env.py::get_aot_dir(), which checkscache_version.startswith(flashinfer_version). The jit-cache wheel has to be rebuilt alongside the tag.Test plan
import flashinferandimport flashinfer.common ROCm with no GPU visible — the cheapest signal, and the one that caught thequantizationpackage defect.pre-commit run -a— clean, without reformatting a single upstream file.pytest -rs --noconftest tests/rocm/test_{upstream_canary,git_describe,gen_arch_support_matrix,build_backend}.py— 142 passed, no GPU.pytest tests/rocm/test_{benchmark_harness,jit_env,comm_import_gate}.pyon gfx942 — 51 passed.v0.6.18tag, restored.pytest -n 4 --reruns 2 -m "not slow" tests/rocm/on gfx942 (MI300X) and gfx950 (MI350X), both from a detachedpin-runworktree atacf125ecd— zero failures on both, exit 0,--checkconfirming neither tree moved mid-run. An earlier pass atfe3c4b4c0returned an identical 11 failures on both arches, all in two host-only files, fixed inf1d69e483. Only7bded9b66(packaging metadata and a CI workflow) landed after the suite SHA; the two test files that parsepyproject.tomlwere re-run against it.validfrom one declaration), the fp16/bf16 cast-back (dropped thecopy_),valid(droppedmark_all_valid), and the import gate (insert→append). Every one fails with the fix reverted — the gate test did not, first time round, which is why it now gates a module that exists.fmha_v2files and nothing else. A scan of both for the licence string comes back empty, and the installed wheel imports with the module absent.scripts/upstream_canary.pyon both refs. Againstv0.6.18:clean merge -- nothing to do, and no upstream changes to any of the 25 forked headers — the port is exactly in sync with the tag it targets. Againstupstream/main: 38 conflicted, 32 of them code, but see below before reading that as fork delta.Follow-ups
v0.6.18is not onupstream/main. It is cut from a release branch that diverged at61a6c6518(2026-08-19), sogit merge-base v0.6.18 upstream/mainis that commit and not the tag. The canary's 38 conflicts againstmaintherefore attribute the whole v0.6.18 release delta to "ours" —fused_moe/runners.py,cute_dsl/tuner.pyand the moe tests dominate the list, and none of them is fork-edited. The honest number for this PR is thev0.6.18one: clean. It does mean the next sync should pick a tag or amaincommit deliberately rather than assume tags are ancestors.mainseven already show churn (attention/prefill.cuh+102/−42,norm.cuh+137/−0, and five smaller). They merge cleanly forever and go stale silently, so the drift report is the only signal.prefill.py(+3019/−195),decode.py(+1514/−437) and turnedmla.pyinto a package.flashinfer/rocm/{decode,prefill,mla}.pyare standalone and unaffected by the merge, but a vLLM or SGLang caller on the 0.6.18 pin will find gaps.block_sparse_indices_to_vector_sparse_offsetsis one concrete instance: upstream removed the Python op, and onlycsrc/rocm/still implements it, with nothing calling it.