Skip to content

Sync with upstream v0.6.18 - #344

Merged
demandal25 merged 1148 commits into
amd-integrationfrom
upgrade/amd-integration+v0.6.18
Sep 2, 2026
Merged

Sync with upstream v0.6.18#344
demandal25 merged 1148 commits into
amd-integrationfrom
upgrade/amd-integration+v0.6.18

Conversation

@demandal25

@demandal25 demandal25 commented Sep 1, 2026

Copy link
Copy Markdown
Collaborator

Summary

Syncs the port from upstream v0.5.3 to v0.6.18 — 1125 upstream commits and 2224 files since the merge base 2628bebcf (2025-11-20). vLLM and SGLang both pin flashinfer-python==0.6.18 exactly 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: ad239a313 is the merge with every shared code file taken from upstream verbatim, and each commit after it restores one subsystem.

What changed

The merge — ad239a313

git merge-tree raised 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, so git diff v0.6.18 over those paths is empty at that commit and the ROCm side comes back in reviewable pieces afterwards.

  • flashinfer/norm.pyflashinfer/norm/ — upstream turned the module into a package, so the deletion is accepted and the dispatch is re-homed in fc730dd08. This is the one conflict that fails silently if fumbled: leaving norm.py beside the package resolves to the package and the ROCm routing just disappears, with no import error.
  • DeletedCHANGELOG.md and tests/utils/test_block_sparse_indices_to_vector_sparse_offsets.py (both gone upstream); version.txt stays deleted, the port versions from setuptools-scm.
  • OursREADME.md, CONTRIBUTING.md, benchmarks/README.md, CLAUDE.md, .claude/skills/benchmark-kernel/SKILL.md, pyproject.toml. .gitignore is the union.
  • pytest.ini is 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.py is unchanged in shape; two things shrank, because upstream deleted get_nvshmem_{include,lib}_dirs outright and now does the extra_include_pathsPath conversion itself. core.py needed real work: upstream refactored JitSpec into an ABC with a JitSpecNvcc subclass and moved build_and_load() onto the base, so the port's old class_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 switched build() to run ninja with -C self.build_dir instead of the JIT root; the ROCm generator still emitted $name/<obj>, which under the new workdir would have nested every artifact one level deeper.
  • The four __init__ files (275f64be8) — mechanical. Upstream's export list re-indented under if IS_CUDA:, the port's elif IS_HIP: / else: arms re-attached unchanged. No symbol the port binds was renamed or dropped.
  • flashinfer/{activation,page,pod,rope}.py (2d1228dca) — the backend="auto"|"native"|"aiter" parameter and its maybe_* short-circuit. pod.py absorbs a fifth CUDA-only plan() argument (uniform_q_len) as one more element of the tail the port already splits off; rope.py grew a seventh _fake_* op.
  • flashinfer/norm/__init__.py (fc730dd08) — plus one addition, below.
  • The benchmark harness (3b742fa08, 4b20e8011) — against a routines/attention.py upstream changed by +1991/−177.

Defects this surfaced

  • import flashinfer was 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, matching comm/__init__.py. packbits is the only member ROCm compiles.
  • The wheel would have been indistinguishable from upstream's (a5fc63be4). scripts/git_describe_rocm.py took the closest ancestor tag regardless of shape; after the merge that is upstream's own v0.6.18 at distance 364 against v0.5.3+amd.2 at 1249, and a bare tag yields a version with no +amd local segment and no error to say so. Fork tags now win at any distance. Restricting the git tag -l glob to v*+amd.* instead would have broken six of the thirteen existing cases.
  • Norm silently stops defaulting to the native kernel if nvidia-cutlass-dsl is ever installed (fc730dd08). Upstream routes norm through CuTe DSL by default and ROCm escaped only via an except ImportError — i.e. because the package happens to be absent. _USE_CUDA_NORM is now forced under IS_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.
  • The benchmark harness aborted before running anything (4b20e8011). Three module-level imports added in 0.6.18 fail on ROCm, so parse_args died on ImportError: cannot import name 'autotune' from 'flashinfer' — which made every ROCm hook restored in 3b742fa08 unreachable. All three uses are CUDA-only paths and move to their call sites.
  • Every sampling call raised at dispatch (fe3c4b4c0). flashinfer/sampling.py is shared verbatim and 0.6.18 widened 10 of its 11 entry points, so the first call died with sampling::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_probs ops gained a valid output, and the renorm ops gained multi-CTA scratch plus an is_deterministic switch for AIR top-p. None maps onto a ROCm kernel — a per-request seed is rejected rather than silently collapsed to seed_val, valid is 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 builds csrc/topk.cu, which csrc/rocm does not have, so it failed inside ninja rather than at import; gated off on HIP. tests/rocm/test_sampling_binding_abi.py compares the two sides statically — no GPU, and it would have caught all three widenings at merge time.
  • tests/attention/test_page.py stopped collecting, so a bare pytest died before running anything (d023ac824). 0.6.18 gave it a module-level nvfp4 import. It is dropped from testpaths — its one HIP-relevant case is covered further by tests/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; a tests/rocm/-only run cannot.
  • The CUDA-only import gate was never installed on ROCm (a8523a350). 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 up only for callers that imported comm first. Every existing test hid it by importing comm itself. Called from flashinfer/rocm/api.py now.
  • The release artifacts carried NVIDIA-licensed sources (7bded9b66). csrc/fmha_v2/ and flashinfer/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 licensed generator_utils.py. The wheel case needed exclude-package-data, not the packages.find exclude: the directory has no __init__.py, so the files arrive as package datainclude-package-data defaults true for a pyproject project and setuptools-scm's file finder then sweeps every tracked file under flashinfer/ in.
  • pre-commit CI was red on the new Validate CUDA configuration step (7bded9b66). Not noise: with pyproject.toml reshaped to suit its hand-rolled TOML parser it then demands nvidia-cutlass-dsl[cu13]>=4.6.2a0 in a cu13 extra, which a ROCm wheel cannot declare. The step is dropped.
  • Two tests were asserting against the wrong file (4b20e8011). test_hip_gqa_group_sizes_match_the_kernel_dispatch read upstream's include/flashinfer/utils.cuh, which gained a group_size == 6 arm; the HIP decode kernel compiles against the forked rocm/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_rocm named gen_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):

  • Ten upstream workflows. Seven fire on this fork's own activity and cannot work here — pr-labeler.yml (pull_request_target, no branch filter, needs FLASHINFER_BOT_TOKEN) would be red on every PR and issue-claim.yml (cron */30) roughly 48 times a day. They are dropped. pr-test.yml, pr-api-doc-checks.yml and pr-api-doc-comment.yml are branches: [main] and stay dormant, so they are left alone — deleting an upstream file buys a permanent conflict on every future sync.
  • pre-commit run -a wanted to rewrite 24 upstream files. The port kept its own pyproject.toml through the merge and so never inherited the [tool.ruff] extend-exclude block 0.6.18 added for generated and vendored kernels; without it, 16 upstream files reformat. markdownlint is fork-added — 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. Both fixes are in the config, not the files.
  • flashinfer.comm.mixed_comm joins CUDA_ONLY_MODULES: it is where the NVSHMEM build landed, and it renders .cu templates into FLASHINFER_GEN_SRC_DIR before reaching import 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 with 128f9dfd5 mid-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 in 61c11fdb8.

Both architectures. The merge changes zero files under include/flashinfer/rocm, csrc/rocm, flashinfer/rocm, tests/rocm, benchmarks/rocm or docs/rocm, and the ROCm jinja templates include only flashinfer/rocm/* headers, so no upstream header reaches the HIP build. The 196-file include/ 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.py derives its base from the fork's own *+amd.* tag, so it self-corrects the moment v0.6.18+amd.1 is cut. Until then it derives v0.5.3, which is not an ancestor of the merge, and reports 827 owned .py files instead of 39 — green, because the arch-caps job only asserts owned is 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 at f58311959: tagging v0.6.18+amd.1 moves the version from v0.5.3+amd.2.dev1266 to v0.6.18+amd.1, the derived base to 69ff11fc4 ("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-cache wheel will hit a RuntimeError from flashinfer/jit/rocm/env.py::get_aot_dir(), which checks cache_version.startswith(flashinfer_version). The jit-cache wheel has to be rebuilt alongside the tag.

Test plan

  • import flashinfer and import flashinfer.comm on ROCm with no GPU visible — the cheapest signal, and the one that caught the quantization package 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}.py on gfx942 — 51 passed.
  • A/B on the version fix: reverted it, confirmed the new case picks the bare v0.6.18 tag, restored.
  • pytest -n 4 --reruns 2 -m "not slow" tests/rocm/ on gfx942 (MI300X) and gfx950 (MI350X), both from a detached pin-run worktree at acf125ecd — zero failures on both, exit 0, --check confirming neither tree moved mid-run. An earlier pass at fe3c4b4c0 returned an identical 11 failures on both arches, all in two host-only files, fixed in f1d69e483. Only 7bded9b66 (packaging metadata and a CI workflow) landed after the suite SHA; the two test files that parse pyproject.toml were re-run against it.
  • A/B on each new test that asserts a behaviour this PR adds: the binding-ABI check (dropped valid from one declaration), the fp16/bf16 cast-back (dropped the copy_), valid (dropped mark_all_valid), and the import gate (insertappend). 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.
  • Wheel and sdist built and their full inventories diffed: 986 -> 982 and 2101 -> 2027 entries, the delta being exactly the NVIDIA-licensed fmha_v2 files 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.py on both refs. Against v0.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. Against upstream/main: 38 conflicted, 32 of them code, but see below before reading that as fork delta.

Follow-ups

  • v0.6.18 is not on upstream/main. It is cut from a release branch that diverged at 61a6c6518 (2026-08-19), so git merge-base v0.6.18 upstream/main is that commit and not the tag. The canary's 38 conflicts against main therefore attribute the whole v0.6.18 release delta to "ours" — fused_moe/runners.py, cute_dsl/tuner.py and the moe tests dominate the list, and none of them is fork-edited. The honest number for this PR is the v0.6.18 one: clean. It does mean the next sync should pick a tag or a main commit deliberately rather than assume tags are ancestors.
  • Forked-header drift is zero right now, and that is the moment it starts accruing. The 25 forked headers match their v0.6.18 counterparts, so the canary reports nothing; against main seven 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.
  • ROCm public-API parity with 0.6.18. Upstream reworked prefill.py (+3019/−195), decode.py (+1514/−437) and turned mla.py into a package. flashinfer/rocm/{decode,prefill,mla}.py are 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_offsets is one concrete instance: upstream removed the Python op, and only csrc/rocm/ still implements it, with nothing calling it.

mhoqueanik and others added 30 commits July 28, 2026 07:04
…, 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>
demandal25 and others added 6 commits August 31, 2026 23:08
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>
Copilot AI lite review requested due to automatic review settings September 1, 2026 13:57

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Comment thread flashinfer/version.py
Comment thread tests/utils/test_jit_warmup.py
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.
Copilot AI review requested due to automatic review settings September 1, 2026 21:43

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 Needs a closer look

Review details
  • Files reviewed: 65/2232 changed files
  • Comments generated: 1
  • Review effort level: Lite

Comment thread csrc/fmha_v2/fmha/hopper/utils_gmma.h
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.
Copilot AI review requested due to automatic review settings September 2, 2026 00:37

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 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
demandal25 merged commit de79634 into amd-integration Sep 2, 2026
3 checks passed
@demandal25
demandal25 deleted the upgrade/amd-integration+v0.6.18 branch September 2, 2026 01:59
demandal25 added a commit that referenced this pull request Sep 2, 2026
Record the v0.6.18 merge that #344's squash discarded
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.