You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
TL;DR. Two public functions named recurrent_kda disagree on their default backend, and the top-level default is not a safe superset of the other. The unification itself is not proposed for 0.7 — it touches released API and belongs on a normal deprecation schedule. Four things do want attention before the cut:
A bug on the default path:#4935 — backend="auto" decode raises instead of falling back, contradicting its own docstring, which says "auto" selects Cake for its native contract and "otherwise preserves CuTe DSL" (recurrent_kda.py:1873-1878). Fixing it is a restoration of the documented contract, not a behaviour change. The only release-critical engineering item here; fix proposed in #5037.
A holding action: six of the ten top-level KDA symbols have never been in a release. 0.7 should not advertise them as stable API. The plan is to mark the unreleased surface experimental, which says plainly that these signatures may change as the unification below lands. See Holding action.
One comment on an open PR:#4779 adds a fourth backend literal to one facade only. The cheapest intervention on this list, and it is open now.
One housekeeping close:#4367 is already fixed on main.
The rest is background and a proposed sequence.
Context
Status as of 2026-09-03 on FlashInfer main (7a3c04f0).
This supersedes #4483 ([KDA] Unify flashinfer.recurrent_kda API, 2026-08-12). That audit is still worth reading as the historical record, but enough has changed in three weeks — a backend= parameter landed, the top-level surface grew from 3 symbols to 10, and a release was cut in the middle of it — that patching it in place would have been more confusing than restating it. Relationship to #4483 summarises what moved.
FlashInfer has two decorated functions named recurrent_kda registering the same trace template. Since #4483, backend=has been added to the top-level facade (#4605), which resolves the original discoverability gap. But the two facades now ship different defaults, and the top-level default is not a safe superset of the other: backend="auto" decode commits to a frozen Cake kernel and raises instead of falling back to CuTe DSL. That defect is filed separately as #4935; this issue covers the public Python API contract only.
The KDA public surface also grew from the 3 top-level symbols #4483 audited to 10 — of which v0.6.18 shipped 4 — adding a training/backward family and a plan-and-run wrapper class that #4483 did not cover.
Six of those ten symbols, and the entire backend= parameter on the top-level facade, have never been in a release — see Release status. That is the most important input to the plan below: the parts of this design that are still wrong can mostly be fixed outright, but only until 0.7 is cut. Since that cut is close, Holding action proposes narrowing the advertised surface back to what v0.6.18 shipped rather than rushing the consolidation.
This is a design / tracking issue. Related CAKE progress is tracked in #4254. The same problem class on the sibling GDN op is already tracked in #2687, with an in-flight implementation in #2706 — that pair is the closest precedent for the plan below and should be aligned with rather than duplicated.
Suggested implementation sequence
Steps are ordered by dependency, but the release window is the scheduling constraint: anything touching unreleased surface is cheap now and locked after 0.7.
Step
Scope
Compatibility
Before 0.7?
Status
0 (done)
Add backend= to flashinfer.kda.recurrent_kda; wire Cake decode and prefill eligibility through it
Comment on #4779 and #4968, asking that each new backend literal cover both facades or go in as an experimental backend under #4880
Two comments; done
Yes
Highest leverage per unit effort: both are open now and each would widen the divergence this issue exists to stop. #4968 adds "cudnn" to kda.py's enum and leaves kda_decode.py untouched — the same shape as #4779's "vibecuda". Raised on #4968 and #4779; #4779 also AOT-registers its backend, which #4880's policy forbids for an experimental one
0c
Housekeeping: close #4367 as already fixed (verified passing on main)
None
no
1
Prerequisite (#5037, open): fix #4935 so the fallback is driven by the resolved selector result rather than any duplicated predicate, with a negative test asserting output and state-pool equality on both sides of the one-warp threshold
Bug fix on unreleased surface
Yes
Blocks step 5. Neither reordering nor a gate-side predicate works — see #4935
2
Document "auto" as the top-level default, keeping the decode facade's released "cute-dsl" default untouched
Free choice while unreleased
Yes
Depends on step 1. No value is safe for a phase-neutral facade until "auto" decode is total — see remedy 2 above
2b
Confirm the "auto" decode kernel swap against v0.6.18's effective "cute-dsl" routing on the intended serving path
Unreleased
Yes
Blocks ratifying "auto" in step 2. Measured: Cake wins at every batch size under CUDA-graph capture; the eager small-batch loss is host dispatch
3
Give the training/backward family real trace templates (or state the omission), and reconsider the _kda_training_impl split
Free while unreleased
no
Problem 8. Disappears entirely if #4965 lands; otherwise is_experimental keeps them out of the stable inventory, so 3b-i removes the urgency
3b-i
Holding action (#5040, open): mark the unreleased top-level symbols experimental under #4880 — RecurrentKDAPrefillWrapper per-method on plan / run, plus the five training/backward symbols only if #4965 has not landed
Non-breaking for callers (warns once)
Yes
No kernel or wheel change. #5040 also relocates the wrapper's own two tests to tests/experimental/, leaving the three CuTe-DSL-numerics tests that merely drive the wrapper in the stable lane
3b-ii
Containment, only if #4965 does not land: relocate kda_backward.py + _kda_training_impl.py behind thin gated entry points, move their tests to tests/experimental/, de-register the ten aot.py sites, file the owner / tracking issue / graduation plan
Wheel contents change; ~91 s first-call JIT; 1,522 test lines leave the default lanes
no — first release after 0.7
Superseded by #4965's removal, which deletes the same code instead of moving it
3c
Settle the SM120a tranche. The __all__ half is done in #5040 (open), which drops the three unreleased kda_kernels.__all__ exports and restores __all__ to byte-identical with v0.6.18; a zero-consumer change, since dispatch reads them as attributes and every doc, test and benchmark uses the deeper kda_kernels.sm120_prefill path. Still open: the explicit backend= literal and gating the automatic branch with experimental_auto_backends_allowed()
Unreleased throughout
Yes for the __all__ half; the gate can follow
Independent of 3b-i. The gate is deliberately not in #5040: it changes "auto" prefill routing on CC 12.0 and mostly by raising (see remedy 2 above), so it wants its own PR validated on a CC 12.0 device
4
Factor a shared _dispatch_recurrent_kda_decode used by both facades; decide which facade owns the recurrent_kda trace
Refactor, no behavior change
no
5
Point kda_decode.recurrent_kda at that dispatcher with omitted→"cute-dsl" + deprecate
Soft deprecation of released API
no
Needs steps 1, 4
6
Give the two recurrent_kda traces distinguishable labels; extend the trace template and retag stage
Additive
no
7
Retarget decode benchmarks, test_recurrent_kda.py, Cake export tests, and fi_trace to the canonical facade; add top-level fi_api trace coverage
Test-only (touches many call sites)
no
8
Deprecate then remove kda_kernels noun aliases (un-shadows submodules); consolidate docs pages; document or internalize FLASHINFER_KDA_T1_FAST_PATH
Soft deprecation of released API
no
The time-critical items are 0b, 1, 2, 2b, 3b-i and 3c; all are confined to main-only surface. Everything else concerns released API and can proceed on a normal deprecation schedule.
Stated plainly: if step 3b-i lands, the release-critical engineering reduces to two items — step 1, and step 3c's disposal of the three unreleased kda_kernels SM120a exports, because the top-level backend= parameter and its default would then be the sole unreleased inference surface that a 0.7 cut freezes — and step 2 cannot be settled without it, since no backend value is safe across both phases until "auto" decode is total. Step 0b is one comment; step 3's subject disappears with 3b-i.
The benchmark in step 2b is not a formality. On a B200 (H=32, K=V=128, T=1, eager, 200 iterations) "auto" decode selects Cake and crosses over against "cute-dsl" at about B = 35: below it "auto" is up to 3.1x slower, above it up to 7.5x faster (0.32x at B=4, 1.90x at B=64, 7.53x at B=256). Cake's flat ~108 µs is host dispatch rather than device time, and under CUDA-graph capture the picture inverts completely: Cake wins at every batch size, by 3.7x to 7.9x. So the eager small-batch loss is ~75 µs of Python-side enqueue that capture erases, and no size condition on the gate is warranted — the default is right for the graph-captured serving path it exists for. What remains worth confirming is that the intended consumers do capture.
If step 1 slips, the contingency is not to ship "auto" as-is. Disabling the speculative "auto" → Cake decode gate is a one-line change to auto_unbounded_softplus_candidate (recurrent_kda.py:2063-2064) that makes "auto" decode identical to "cute-dsl" — bit-exact on every trigger class with a non-zero pool and non-identity indices, backend="cake" strictness intact, and exactly the routing a defaulted v0.6.18 call took. It costs the Cake decode win at serving batch sizes for one release and requires retargeting six parametrizations of test_t1_unbounded_softplus_auto_route_* to backend="cake". Land step 1; keep this in reserve.
Current public surface (main)
Full inventory: both facades, the ten top-level symbols, and the docs split
flashinfer/kda_training.py (re-export of private _kda_training_impl.py)
bare, no trace / n/a
paired training API
New. Own paired-function convention
Top-level __init__.py re-exports 10 KDA symbols (__init__.py:165-184): recurrent_kda and RecurrentKDAPrefillWrapper from kda; fused_kda_decode and packed_kda_decode from kda_decode; RecurrentKDAPrefillWorkspace from kda_prefill; recurrent_kda_backward and RecurrentKDABackwardWorkspace from kda_backward; and recurrent_kda_training_forward, recurrent_kda_training_backward, RecurrentKDATrainingContext from kda_training.
flashinfer.kda is reachable as an attribute (the package-root from .kda import … binds the submodule) but is not an intentional re-export.
flashinfer.kda_kernels still re-exports three noun aliases (recurrent_kda, fused_kda_decode, and now packed_kda_decode) alongside the run_* backends. Its __all__ has grown from 6 entries to 9 since v0.6.18; all three additions are the SM120a prefill helpers, which are unreleased. Two of the three are not merely redundant — they shadow a submodule of the same name, so attribute access cannot reach the module:
K.recurrent_kda -> function (shadows kda_kernels/recurrent_kda.py)
K.fused_kda_decode -> function (shadows kda_kernels/fused_kda_decode.py)
K.packed_kda_decode -> function (no same-named module; impl is cake_packed_kda_decode.py)
importlib.import_module("flashinfer.kda_kernels.recurrent_kda") -> module
In-repo consumers that need the shadowed modules work around this with importlib.import_module. That makes the alias cleanup a correctness fix for the first two, not cosmetics. Separately, kda_kernels/__init__.py does now carry an explicit note that packed_kda_decode_cute is an internal implementation module — the desired precedent, applied inconsistently with the shadowing noun aliases directly above it.
Docs are split across five pages: docs/api/kda.rst, kda_decode.rst, kda_prefill.rst, and the new kda_backward.rst and kda_training.rst (both deleted by #4965). kda_decode.rst still autosummarizes recurrent_kda, packed_kda_decode, and fused_kda_decode. kda_prefill.rst documents RecurrentKDAPrefillWrapper, which is defined in kda.py.
In-repo usage today: Prefill tests/benchmarks use the top-level / kda facade (bench_recurrent_kda_prefill.py, tests/kda/test_recurrent_kda_prefill.py). Decode paths — including Cake export tests and the main CuTe DSL correctness suite tests/kda/test_recurrent_kda.py, plus bench_recurrent_kda.py, bench_recurrent_kda_decode_export.py, and bench_recurrent_kda_decode_split_sweep.py — import flashinfer.kda_decode.recurrent_kda. The newer tests/kda/test_packed_kda_decode_cute.py uses top-level flashinfer.recurrent_kda, so the codebase now exercises both facades for decode-shaped calls. fi_trace fixtures still pin only the decode facade.
The split is also why CI does not catch #4935, though the mechanism is subtler than "the top-level default is untested". Across tests/kda/, the "auto" gate matches on 6 of 190 run_recurrent_kda calls, all six with an explicitbackend="auto" and all from one test that asserts Cake is selected — so it cannot detect a missing fallback. The three defaulted top-level calls in test_packed_kda_decode_cute.py all miss the gate, two on lower_bound=-5.0 and the third on use_gate_in_kernel. The gap is a missing negative test, not missing coverage, which makes it cheap to close.
Release status: what is actually shipped
Most of the surface this issue describes has never been in a release. That widens the design freedom considerably, and it also puts a clock on it.
v0.6.18 (published 2026-08-29, tag commit 69ff11fc) is the current stable release. It is not an ancestor of main: the release branch was cut from main at 61a6c651 on 2026-08-19 (author date 01:03Z; 08-18 in US-Pacific, which is how it appears in git log locally), then took cherry-picks through 08-28. So "did it ship?" cannot be inferred from merge dates — the classification below comes from reading the v0.6.18 tree directly.
So v0.6.18 exports 4 KDA symbols where main exports 10.
No other channel carries the newer surface either. PyPI serves flashinfer-python0.6.18 as latest (plus a 0.6.18rc10 prerelease, checked and clean); the release-v0.6.18 branch head is the tag itself; there are no 0.6.18.post* releases, which is informative rather than accidental since v0.6.16 had four. Nightly tagging stops at nightly-v0.6.18-20260819 — one day before #4605 — and nightlies track main rather than the release branch, which is what makes that check meaningful. No nightly contains a top-level backend= or the auto_unbounded_softplus_candidate gate.
Per-item breakdown: what is changeable outright vs. what carries a compatibility constraint
Unreleased — changeable outright, no deprecation cycle
The backend= parameter on the top-level facade, in its entirety. Released flashinfer.recurrent_kda has no such parameter; the only backend= in the released kda.py is an internal hardcoded backend="cake" prefill call. "auto" as the top-level default is therefore unreleased.
RecurrentKDAPrefillWrapper and the entire training/backward family (five symbols), with their modules kda_backward.py, kda_training.py, _kda_training_impl.py, _kda_training_dispatch.py, kda_prefill_cute.py and docs pages kda_backward.rst, kda_training.rst. The bare-decorator/no-trace choice and the private _*_impl split (problem 8) can simply be corrected. The unreleased footprint is wider than the public symbols suggest: 15 KDA-related Python modules exist on main and not in v0.6.18, the rest being kda_kernels/kda_chunked_bt16.py, the four kda_kernels/sm120_prefill/ files, and five jit/ modules (_kda_jit_common.py, cake_kda.py, cake_kda_decode.py, flash_kda_backward.py, flash_kda_training.py).
The whole SM120a prefill tranche.kda_kernels.__all__ grew from 6 entries to 9, and all three additions are SM120a: can_implement_kda_prefill_sm120, clear_kda_prefill_sm120_caches, run_kda_prefill_sm120. Behind them sits a new flashinfer/kda_kernels/sm120_prefill/ subpackage (decomp.py, fused.py, runtime.py), the _sm120_kda_prefill_* eligibility and rejection machinery in kda_prefill.py, RecurrentKDAPrefillWorkspace._sm120_state, and the "auto" prefill route that tries SM120a first (kda.py:302-306). None of it is in v0.6.18.
Released — real compatibility constraints
Note that the core complaint of this issue is in the released set, so consolidation still needs a deprecation story:
Both recurrent_kda facades shipped under the same name, and both were already decorated with recurrent_kda_trace — the duplicate name and the trace-registry collision (problems 1 and 2) are released.
Problem 3(b) is released: _is_plain_multi_token_prefill and the top-level facade's use of it both shipped, while the decode facade never classified phase.
kda_decode.recurrent_kda's keyword-only backend: Literal["cute-dsl", "cake"] = "cute-dsl". main has since widened that enum to include "auto" (feat(cake_kda): add native unbounded-softplus Kimi-Linear kernels #4535), which is purely additive; changing the default would be a breaking behavior change.
seq_order, prefill_workspace, and the checkpoint kwargs on the top-level facade.
The kda_kernels noun aliases and their submodule shadowing (shim item 4).
FLASHINFER_KDA_T1_FAST_PATH (problem 6).
The three-page docs split.
Problem 3(a) — the default divergence — is the one part of problem 3 that is unreleased, precisely because the released top-level facade had no backend at all.
Timing
main is still stamped 0.6.18, and no 0.7 branch or rc tag exists yet. Recent stable cadence is roughly two to three weeks (v0.6.17 2026-08-11, v0.6.18 2026-08-29), so the next cut is plausibly mid-September. Once 0.7 ships backend="auto" as the top-level default, changing it presumably owes users a deprecation cycle rather than being a straight correction, and the cheap fixes here turn into expensive ones. The unreleased-surface decisions should land before the next release cut.
Holding action: what 0.7 should advertise
The plan is not to advertise the new KDA surface as stable in 0.7. Six of the ten top-level symbols have never shipped, and several are ones this issue proposes to change. Marking them experimental states that directly: they remain usable, but their signatures may change as the unification below lands, and nobody acquires a compatibility claim on them in the meantime.
Update (2026-09-04): five of those six are now proposed for removal rather than marking. Per @YangXu1990uiuc's comment, FlashInfer and cuDNN Frontend are splitting KDA by phase — inference here, training there, where cudnn.linear_attention.ops.kimi_delta_attention already provides fwd+bwd with autograd and the CAKE training kernels will land as an engine under it. #4965 implements the removal of recurrent_kda_training_forward, recurrent_kda_training_backward, RecurrentKDATrainingContext, recurrent_kda_backward and RecurrentKDABackwardWorkspace. Since none has shipped and a GitHub-wide search finds no external caller, removal before the cut needs no deprecation cycle, and it disposes of problem 8 and the aot.py registrations outright instead of relocating them. That leaves RecurrentKDAPrefillWrapper as the only unreleased top-level symbol needing the experimental treatment. #4965 is still a draft pending its cuDNN Frontend counterpart (NVIDIA/cudnn-frontend#912), so the marking below remains the fallback for the training family if the cut lands first. The rest of this section is written for that fallback; ordinary prefill, including the m128_n16_checkpoint variant and the checkpoint kwargs that serve SGLang prefix caching, is untouched either way.
The mechanism is #4880, which merged on 2026-09-04. Marking a symbol is non-breaking: @flashinfer_experimental_api emits an ExperimentalWarning once rather than raising, and its docs are explicit that calling an experimental API is itself the opt-in, requiring no environment variable. What the env var FLASHINFER_ALLOW_EXPERIMENTAL_AUTO_BACKENDS=1 gates is automatic selection — whether backend="auto" may pick a backend marked @experimental_backend. trace is optional for experimental APIs and is_experimental excludes them from the stable trace inventory, which incidentally sanctions the training family's bare-decorator state (problem 8) rather than leaving it an inconsistency.
Applying it means following #4880's own guidelines rather than bolting a decorator onto the current layout. Concretely, for the KDA surface:
Decorate methods, not classes.@flashinfer_experimental_api returns a function, so applying it to a class silently replaces the class with a function and breaks isinstance, subclassing and Sphinx autoclass — and docs/api/kda_prefill.rst autosummarizes RecurrentKDAPrefillWrapper, so that would be a visible docs regression. Per-method decoration avoids this and is already the codebase's convention for class-based public ops: 98 indented @flashinfer_api sites, including sparse.py:1224 where BlockSparseAttentionWrapper.run carries a full trace template, against zero applications to a class. That covers all six symbols:
remove (#4965); as fallback, transitively covered — inert without recurrent_kda_backward
RecurrentKDATrainingContext
class
remove (#4965); as fallback, transitively covered — a returned tape dataclass, never user-constructed
Worth adding an isinstance(f, type) guard to the decorator so the degradation is an error rather than silent; raised on #4880. Note that per-method decoration warns but does not document: RecurrentKDABackwardWorkspace and RecurrentKDATrainingContext have their own autosummary entries and would render as ordinary stable pages, so each needs an explicit .. warning:: in its class docstring.
Marking and containment are separate pieces of work, and only marking belongs before the cut. Decorating the symbols is hours of work with no kernel, wheel or CI-lane change. Containment is not: #4880's CONTRIBUTING.md requires experimental tests to live in tests/experimental/, and its pytest.ini sets norecursedirs = test_helpers experimental, so complying moves tests/kda/test_recurrent_kda_backward.py and test_recurrent_kda_training.py — 1,522 lines of Blackwell coverage — out of the default lanes days before a cut. De-registering also breaks test_aot_registers_flash_kda_backward, which pins the spec list exactly (tests/jit/test_flash_kda_backward_jit.py:206-213), and the real footprint is ten aot.py sites plus three tests, not four entries: imports at :98-99, capability probes at :553-557, the four appends at :674-679, and detect_sm_capabilities at :1235-1242. The relocation is also not a file move — the containment rule wants a thin core entry point, and flashinfer/kda_training.py is currently a bare re-export, so the 15-parameter signature has to be restated in core, where tests/kda/test_recurrent_kda_training.py:86-121 pins it verbatim and a *args, **kwargs handoff fails. First-call JIT for training users is roughly 91 s per cache directory, measured. None of that is an argument against containment — it is the agreed destination — but nothing in the 0.7contract depends on file placement, so it should land in the first release after. RecurrentKDAPrefillWrapper is exempt either way: it is host-side planning that ends in a call to the stable recurrent_kda facade, so it has no backend implementation to contain.
Relocate and de-register, per the containment rule — the fallback path only, superseded if #4965 lands.#4880's README is explicit that a core entry point may include only the signature, shared validation, the gate check, backend selection and a handoff, and that "if the API is experimental, its implementation lives here too, regardless of the implementation's maturity — this keeps removal a one-directory deletion". So kda_backward.py and _kda_training_impl.py (~3,000 lines) move under flashinfer/experimental/ behind thin gated entry points, and the ten aot.py sites enumerated above come out, leaving the training and backward kernels JIT-only as the policy requires. Two consequences to expect rather than to debate: wheel contents change, and training users take first-call JIT latency. Both are the normal price of the experimental track, and both reverse cleanly when these APIs graduate.
One genuine convergence, worth aligning deliberately.FLASHINFER_ALLOW_EXPERIMENTAL_AUTO_BACKENDS exists precisely so that backend="auto" declines to select experimental backends — which is exactly the shape of the SM120a prefill question. The unreleased surface splits by architecture, and the two halves want different halves of the policy. The six top-level symbols are Blackwell-datacentre only: recurrent_kda_backward and the training trio gate on _SUPPORTED_COMPUTE_CAPABILITIES = ((10, 0), (10, 3)) (kda_backward.py:35, _kda_training_impl.py:22), and RecurrentKDAPrefillWrapper is documented "Compute capability 10.0 and 10.3 only". Those are API surface. The SM120a prefill tranche is backend surface reached through a stable API — the shape @experimental_backend is built for. It cannot use that decorator yet: @experimental_backend marks a @backend_requirement checker, and KDA has none, so it only becomes available after the migration in remedy 2. Until then the hand-rolled equivalent is warn_experimental_backend_once on an explicit branch plus experimental_auto_backends_allowed() guarding the automatic one at kda.py:302 — notrequire_experimental_auto_backends, which raises rather than falling through. Gating SM120a also changes "auto" prefill routing on CC 12.0, and mostly by raising. Of three prefill spellings that work on CC 10.0, only the packed one still succeeds once SM120a is gated off: a fixed-T call (q=[B, T>1, H, 128], no cu_seqlens) falls through to the generic path and then raises Decode only supports T=1 without cu_seqlens, and a call passing prefill_workspace= raises ValueError because SM120a is the only CC-12.0 backend that accepts a workspace (prefill_workspace at kda_prefill.py:3627, the raise at kda.py:473). Neither is a regression against v0.6.18, which had no CC-12.0 prefill backend at all, but the second contradicts a "falls through rather than raises" reading. Worth noting #4880's stated default here is to raise with a message naming the variable, so a fall-through is a deliberate divergence. There is no SM121 question: the gate is exactly _SM120_KDA_SUPPORTED_COMPUTE_CAPABILITIES = {(12, 0)} (kda_prefill.py:3594, mirrored by kda.py:325), so CC 12.1 is served by no KDA prefill backend, released or not.
The architecture predicate is enumerated in eight Python places, four of which independently spell the same pair.((10, 0), (10, 3)) appears verbatim at kda_backward.py:35, _kda_training_impl.py:22, kda_prefill_cute.py:34 and kda_prefill.py:43 — the last two being the prefill entry-point gates. Alongside them: _FLASH_KDA_DECODE_ARCH_BY_COMPUTE_CAPABILITY (recurrent_kda.py:70-75) for decode, and (12, 0) three times over for SM120a — _SM120_KDA_SUPPORTED_COMPUTE_CAPABILITIES (kda_prefill.py:3594), SM120_CAPABILITY (kda_kernels/sm120_prefill/runtime.py:138) and DEVICE_CC (sm120_prefill/fused.py:99). #4710 is a fair illustration rather than a counterexample: it widened the four csrc/kda/*.cuh runtime guards to admit CC 10.7 and refreshed the JIT cache keys that edit invalidates, and its scope note states plainly that the Python entry points still restrict KDA to 10.0/10.3 so end-to-end Rubin support is separate work. The observation is only that the separate work is spread across five sites, and that while "auto" raises on an unmapped capability (#4935) they have to move in lockstep with the kernels.
Marking does not fix #4935."auto" is the top-level default and the defect is in its decode routing, so that needs the change in step 1 regardless.
The marking and the bug fix are independent, and only the bug fix is release-critical. Marking is worth doing before the cut because it is what keeps the unification below cheap, but nothing breaks if it slips. #4935 is different: it is the behaviour of the default path, so shipping it means shipping a broken default.
How we got here
Timeline of the merged PRs that produced the current shape
Added packed_kda_decode as a public noun-form facade (trace + top-level export + docs), superseding #4378. Also added checkpoint kwargs to the top-level facade
Added backend="auto"|"cute-dsl"|"cake" to flashinfer.kda.recurrent_kda and wired CuTe DSL recurrent prefill. This is Step 1 of #4483's proposed sequence
Made "auto" fall back to Cake when CuTe DSL predates cutlass.experimental (<4.7), and made explicit "cute-dsl" raise ImportError there — introduced version-dependent routing. The strict path is currently red in CI (#4911)
Added "auto" to kda_decode.recurrent_kda and run_recurrent_kda, selecting Cake for the equal-head/D128/T=1 unbounded-softplus contract. Source of #4935
Cake serving-native packed decode; superseded by #4445
The v0.6.18 branch cut falls between #4445 (08-16, released) and #4605 (08-20, unreleased). Everything from #4605 onward in that table — the top-level backend= parameter, both CuTe DSL prefill backends, the version gate, and "auto" decode — is main-only.
Concrete problems
Duplicate canonical name. Two @flashinfer_api wrappers share the name recurrent_kda and the same trace template, and are not the same object (flashinfer.kda.recurrent_kda is flashinfer.kda_decode.recurrent_kda → False). The top-level facade still bypasses the decode facade by calling raw _run_recurrent_kda, so decode through the top-level path is never seen by the decode wrapper's own logging/tracing. (It does now forward backend=, so the original "never sets backend" finding no longer applies.)
Trace registry collision. Both facades register in _TRACE_REGISTRY under the label 'recurrent_kda'. The count depends on import scope — plain import flashinfer shows two colliding labels, and importing every submodule shows four:
Two differently-named functions (…_fused_moe, …_fused_moe_nvfp4) in one module sharing a template
block_sparse_run
Two differently-named functions in different modules sharing a label
Only recurrent_kda is two same-named public facades; the other three are deliberate template reuse that merely produces a non-unique label. Any registry-uniqueness fix should distinguish these cases rather than treat them alike.
No test catches this. The registry's duplicate key is (func.__module__, func.__qualname__, label) (tests/trace/template_registry.py:100-103), and the two facades differ in __module__, so the shared label never trips the guard. The collision is not a failing test — it is unpoliced, which is the weaker-looking but more durable argument for fixing it.
recurrent_kda_trace has since gained ssm_state_indices, state_checkpoints, checkpoint_cu_starts, and checkpoint_every_n_tokens, but still does not model backend, seq_order, or prefill_workspace, and is still tagged stage:decode despite dispatching prefill. fi_trace fixtures still pin fi_api:flashinfer.kda_decode.recurrent_kda only, so top-level fi_api strings remain untested.
The two facades disagree on the default, and the top-level default is not a superset. Both delegate decode to the same run_recurrent_kda, so for a fixedbackend value the decode kwargs are identical. Two things still diverge:
(a) The default value. Verified by intercepting the forwarded value:
Because "auto" commits to Cake without falling back (#4935), the top-level default is not a superset of the decode facade's — and the two accepted sets are in fact incomparable, since one guard is skipped under "auto" that "cute-dsl" enforces. The same call can succeed through one import path and raise through the other, in both directions. This is the strongest argument for consolidation, but it is a defect to fix in [Bug] recurrent_kdabackend="auto" decode commits to Cake and raises instead of falling back to CuTe DSL #4935, not something to design around here.
(b) Phase classification runs before delegation.flashinfer.kda.recurrent_kda calls _is_plain_multi_token_prefill first, which classifies on shape only (kda_prefill.py:315-327): no num_spec_tokens, and either fixed T > 1 or packed total_tokens > num_sequences. There is no head-dim, dtype, or architecture guard. So a packed multi-token decode call is claimed by the prefill branch:
packed q=[1,4,H,128], cu_seqlens=[0,2,4], num_spec_tokens=None, precomputed gate:
kda_decode.recurrent_kda (today) -> OK (reaches decode)
kda.recurrent_kda backend='cute-dsl' -> ValueError: does not support this prefill contract
kda.recurrent_kda backend='cake' -> ValueError: does not support this prefill contract
kda.recurrent_kda backend='auto' -> OK (reaches decode, prefill backends ineligible)
The decode facade never runs this classifier. The "auto" outcome here is gate-mode dependent — with the raw Kimi-Linear softplus gate (use_gate_in_kernel=True, lower_bound=None) the same call raises instead, because it then matches the "auto" Cake gate discussed in [Bug] recurrent_kdabackend="auto" decode commits to Cake and raises instead of falling back to CuTe DSL #4935. Either way the conclusion holds and in fact strengthens: no single backend value reproduces the decode facade's current behavior. This is the constraint that shapes the shim design below.
Docs / traces lag the facade. Five Sphinx pages now describe pieces of one recurrent API; kda_decode.rst still autosummarizes recurrent_kda; kda_prefill.rst documents a class defined in kda.py. Both the kda_decode.py docstring and docs/api/kda_decode.rst describe an "auto" → CuTe-DSL decode fallback that does not exist (see [Bug] recurrent_kdabackend="auto" decode commits to Cake and raises instead of falling back to CuTe DSL #4935). kda_decode.rst documents the Cake export set thoroughly — the exported bodies, gate modes, per-T value-row splits and per-architecture split policies — but not the layout boundary: whether a given call is accepted depends on a 259-line selector over strides, pointer alignment and int32-overflow products that no user can predict from the prose. Once "auto" falls back correctly this is a performance-predictability problem rather than a correctness one, but it should be stated somewhere. The trace template gaps and stage:decode tag are covered in problem 2.
An environment variable participates in public dispatch, documented only for contributors.FLASHINFER_KDA_T1_FAST_PATH (default "1", recurrent_kda.py:139) selects a different T=1 decode route. It appears in CLAUDE.md:607 but nowhere in docs/api/, so users of the public API cannot discover it. Which kernel a given public backend= value resolves to therefore depends on process environment. Any unified contract needs to either document this as public API or make it internal-only.
Fragmentation is now happening, not hypothetical. Open PRs are each adding surface in a different style, concurrently editing the same two facades:
Adds flashinfer/kda_vibecuda.pyand widens the top-level enum to Literal["auto","cute-dsl","cake","vibecuda"] on kda.py only — a fourth backend literal that the decode facade will not accept, pushing the two signatures further apart
Exports generated Blackwell recurrent-KDA specializations, fp32 state and long context; edits kda.pyandkda_prefill.py and adds a jit/flash_kda_nvrtc.py compilation path
Worth citing as the counterexample: #4709 adds disable_state_update with backend-specific raises but updates bothkda.py and kda_decode.py, keeping the two signatures in step. That is the behavior this issue wants to make the default expectation.
New surface outside the original scope.kda_backward.py and kda_training.py introduced a paired-function training family, and RecurrentKDAPrefillWrapper introduced a plan-and-run class convention. Neither is covered by the "noun-form facade vs internal run_*" rule below, so that rule needs cases for stateful plan/run wrappers and paired forward/backward ops. Two further inconsistencies came along with them:
All three new ops use bare @flashinfer_api with no trace template, unlike every recurrent_kda/*_kda_decode facade. So "public op" currently means "decorated" in two incompatible ways, and the training family is invisible to fi_trace.
flashinfer/kda_training.py is a 23-line re-export of private flashinfer/_kda_training_impl.py — a module-layout convention used nowhere else in KDA (kda_backward.py keeps its implementation inline). There is a stated reason: the module also defines _get_training_module, an "internal indirection seam used by the training API tests". That justifies the seam; whether it also justifies the split is worth a sentence from the author. Either way it is cosmetic next to the missing trace templates alongside it.
Adjacent: #4896 (production KDA prefill for SGLang prefix caching).
Non-goals
Do not merge fused_kda_decode or packed_kda_decode into recurrent_kda — different serving contracts (packed QKV, conv, RMSNorm; different return types).
Do not fold the training/backward family into the inference facade.
Do not require open kernel PRs to refactor all of KDA; they should either follow the target facade pattern or stay explicitly internal.
Do not change numerical contracts of existing CuTe / Cake kernels in the unification PR.
Current backend semantics, as actually implemented on main:
backend
Behavior
"auto" (default)
Prefill: SM120a CuTe → SM100-family CuTe DSL BT=16 → frozen Cake fallback (Cake also used when nvidia-cutlass-dsl<4.7). Decode: commits to Cake when the equal-head/D128/T=1 unbounded-softplus gate matches, raising rather than falling back if Cake cannot serve it (#4935)
"cute-dsl"
CuTe DSL strictly; no frozen-prefill dispatch and no Cake decode. Raises ValueError on an unsupported plain-prefill contract, ImportError when the DSL predates 4.7
"cake"
Frozen Cake strictly; raises on unsupported prefill contracts and on unexported decode shapes. No silent CuTe fallback
Target semantics: "auto" becomes a true superset — never raising where "cute-dsl" or "cake" would have succeeded. The plan below assumes that fix has landed.
Because none of this parameter has shipped, the fix should be the structural one described in #4935 — resolve backend eligibility before state normalization, so one resolved decision drives both the state convention and the launch — rather than the cheaper "keep the speculative gate and hand-synchronise a second predicate" variant. There is no released behavior to preserve, so there is no reason to accept two predicates that must be kept in step by hand; that fragility is what produced the defect.
The remaining work is not adding this parameter — it is reconciling the second facade against it. Keep seq_order / prefill_workspace / checkpoint kwargs on this facade only (all three are released there).
Compatibility shims
flashinfer.kda_decode.recurrent_kda cannot simply forward to flashinfer.kda.recurrent_kda. No single backend value preserves today's behavior, because the two facades differ in both the default and the phase classifier (problem 3):
So the shim must bypass phase classification and target a shared decode dispatcher directly, rather than re-entering the phase-neutral facade:
Factor the decode tail of kda.recurrent_kda into one internal _dispatch_recurrent_kda_decode(...), used by both facades, giving decode a single contract regardless of entry point. Note this does not by itself fix tracing: an internal dispatcher is undecorated by construction, so problem 1 additionally requires deciding which facade owns the recurrent_kda trace (see the docs/traces section).
kda_decode.recurrent_kda forwards to that dispatcher with omitted backend → "cute-dsl", preserving today's kernel selection exactly.
Accept and forward "cake" and "auto" unchanged.
Emit DeprecationWarning pointing at flashinfer.recurrent_kda (soft; keep for ≥1 minor release).
Do not accept seq_order / prefill_workspace on the shim — those are prefill-only and the shim deliberately does not reach prefill.
Decide the top-level default explicitly, before 0.7. This is a policy choice separate from [Bug] recurrent_kdabackend="auto" decode commits to Cake and raises instead of falling back to CuTe DSL #4935: even once "auto" stops raising, it selects a different decode kernel than "cute-dsl" would. Today, no single value is a safe default for a phase-neutral facade — and that is the sharpest evidence for this whole issue. The two phases have opposite totality properties:
accepts classes 1-3, and accepts class 4's padded layout only below the one-warp threshold; separately raises on a non-contiguous initial_state pool that "auto" accepts
Neither value is total on either phase, and on decode the two accepted sets are incomparable rather than nested (see #4935, Non-comparability); "cake" is strictly narrower than both. The prefill raise says so in its own message: "backend='cute-dsl' requires nvidia-cutlass-dsl>=4.7.0 (cutlass.experimental); backend='auto' falls back to Cake." So defaulting the top-level facade to "cute-dsl" would break defaulted prefill calls on older CuTe DSL, and defaulting it to "auto" ships the known-broken decode default. A single backend knob on a phase-neutral facade is being asked to encode two independent phase-specific decisions, which is problem 3(b) restated as a concrete release blocker.
That resolves the choice rather than blocking it:
"auto" is the right top-level default, but only once step 1 lands. It is already the only value with fallback semantics on both phases; step 1 makes those semantics true on decode as well as prefill. Without step 1 there is no safe default at all, which is the strongest reason to treat it as release-critical rather than merely desirable.
The top-level default is unreleased, so it can be set either way with no deprecation — but note that in v0.6.18 the top-level facade had no backend at all, so a defaulted top-level decode call reached run_recurrent_kda's own "cute-dsl" default. That call has already changed kernel between the release and main — and on the inputs in #4935 it now raises where the v0.6.18 spelling of the same call worked, which makes step 1 a regression fix against the last release rather than merely a correction to unreleased surface. Ratifying "auto" is a silent kernel swap for existing top-level callers, not a neutral blessing of the status quo, and it should be benchmarked rather than assumed.
Adopt @backend_requirement rather than hand-rolling the dispatch.utils.py:1185 already provides the repo's standard mechanism: per-backend eligibility checkers plus an "auto" heuristic, resolved by computing the set of suitable backends first — filtered on compute capability and problem size — and raising only when that set is empty. It is applied at 44 sites across 15 modules (gemm/, grouped_mm/, fused_moe/, quantization/, deep_gemm.py) and at zero sites in KDA, and it composes with @flashinfer_api and a trace template at 43 of those 44 (e.g. attn_scores.py:1020-1021). Its resolve-then-dispatch ordering is what makes commit-then-check hard to write by accident, which is why it is the right target shape rather than a bespoke selector, and it gives the phase asymmetry above a natural home as two decorated entry points with their own checker sets.
Two limits to plan around, though, which is why this is post-0.7 and not a wholesale replacement. Checkers run before normalization (utils.py:1412-1459 validates and then calls; there is no post-normalization hook), so they can host KDA's contract predicates — trigger classes 1-3 — but not its layout predicates, which are evaluated against re-viewed and synthesized tensors that do not exist until :2263-2337 and which produce trigger class 4. The decorator also injects a public skip_check kwarg whose True path bypasses all checking. So the realistic shape is contract predicates in checkers and a late fallback for layout, not one mechanism for both.
The decode facade's "cute-dsl" default is released and is decode-only, so no prefill concern touches it and it should not be changed at all. Its enum already accepts "auto" (kda_decode.py:79, added by feat(cake_kda): add native unbounded-softplus Kimi-Linear kernels #4535), so nothing needs adding there. The two facades converge on "auto" only after step 1 — and if step 1 slips, they should not be converged at this cut.
flashinfer.kda_kernels.run_*: document as backends only; remove the three noun aliases (recurrent_kda, fused_kda_decode, packed_kda_decode) from kda_kernels.__all__. This un-shadows the same-named submodules and lets in-repo consumers drop their importlib.import_module workarounds. Extend the existing packed_kda_decode_cute "not public API" note into a module-wide statement. These aliases shipped in v0.6.18, so the removal needs the same soft-deprecation window as item 1 even though it is a correctness fix.
Docs / traces
Single "start here" page for recurrent KDA (docs/api/kda.rst); kda_decode.rst focuses on fused_kda_decode + packed_kda_decode plus a compatibility note for the old recurrent_kda import; kda_prefill.rst remains workspace/contract detail; consider moving RecurrentKDAPrefillWrapper onto the page for its defining module (Sphinx resolves it correctly today via .. currentmodule:: flashinfer.kda, so this is page organisation rather than breakage); kda_backward.rst / kda_training.rst stay as the training pair.
Document or internalize FLASHINFER_KDA_T1_FAST_PATH.
Extend recurrent_kda_trace for backend, seq_order, and prefill_workspace, retag the stage now that the template covers prefill, and give the two facades distinguishable registry labels (or collapse them once the shim lands). Treat the gdn_decode / cute_dsl_fused_moe / block_sparse_run collisions as a separate, lower-priority cleanup — they are template reuse, not duplicate facades.
Add fi_trace coverage for the top-level facade's fi_api strings, then point decode benches/tests (CuTe + Cake) at the canonical facade.
Rule for new kernels (immediate guidance)
New KDA work should do one of:
Public op: noun-form facade (@flashinfer_apiwith a trace template + top-level/docs export); run_* stays in kda_kernels as backend. State explicitly if a trace template is deliberately omitted, rather than defaulting to a bare decorator as the training ops do.
Public stateful op: plan-and-run class next to its facade, documented in its defining module's page (as RecurrentKDAPrefillWrapper does), or a paired forward/backward module (as kda_training). Keep the implementation in the public module unless there is a stated reason for a private _*_impl split.
Internal/experimental: do not add to public __all__ / top-level; state that clearly in the PR — the precedent set by feat(kda): add packed-input CuTe decode kernel #4417. With #4880 landing in 0.7 this becomes a formal track: @flashinfer_experimental_api for an unstable public API, and flashinfer/experimental/ plus @experimental_backend for an opt-in backend reached through an explicit backend= literal. Prefer that over widening a stable enum.
Do not add a new top-level *_kda_decode noun without stating why it cannot be a backend= value or a kwarg on an existing facade.
Do not introduce a further convention such as "only kda_kernels.run_foo / kda_kernels.<submod>.run_foo is the product API."
Why this is backwards compatible
Compatibility argument, item by item
The compatibility baseline is v0.6.18, not current main.
Existing flashinfer.kda_decode.recurrent_kda(...) calls keep working: the shim targets a decode dispatcher directly, so omitted backend stays "cute-dsl" and the prefill classifier never intercepts a packed multi-token call. Its released two-value backend enum is only ever widened, never narrowed, and its default is never changed.
Released fused_kda_decode / packed_kda_decode / RecurrentKDAPrefillWorkspace, and the released top-level seq_order / prefill_workspace / checkpoint kwargs, are unchanged.
The kda_kernels noun aliases and FLASHINFER_KDA_T1_FAST_PATH shipped, so both get a deprecation window rather than removal in place.
No forced migration in the first release; warnings guide callers toward one import path.
The training/backward family and RecurrentKDAPrefillWrapper are deliberately not held compatible with main — they are unreleased, so step 3 revises them freely. "Backwards compatible" makes no claim about main-only surface.
#4483 — the 2026-08-12 audit this supersedes; closed 2026-09-03 as completed, "in favor of updated issue"
#4935 — "auto" decode commits to Cake and raises instead of falling back
#5042 — found while reviewing #5037: backend="cute-dsl" dense decode gathers and scatters the state pool in Python, so a negative ssm_state_indices entry wraps around and overwrites the last slot. Released in v0.6.18, silent, and independent of the unification plan — listed here only so the KDA decode defects stay in one place
#4965 — draft removal of the unreleased training/backward family, per the FlashInfer/cuDNN phase split; pairs with NVIDIA/cudnn-frontend#912
#4968 — adds a "cudnn" backend literal to kda.py only; step 0b
#4880 — experimental API/backend policy, @flashinfer_experimental_api, flashinfer.experimental; landing in 0.7; the mechanism for the holding action above
#2687 — [RFC] Unified GDN Decode/Prefill API; the same problem class on the sibling op
#2706 — feat(gdn): add unified decode API and deprecation shims; in-flight implementation of this exact pattern for GDN
#4367 — filed against v0.6.17 for tests/trace/test_template_registry.py::test_registration_module_inventory_is_complete, asserting Extra items in the right set: 'flashinfer.kda'. Already fixed and should be closed: flashinfer.kda is now in the static inventory (tests/trace/template_registry.py:74) and the test passes on main (verified). It is evidence that introducing the second facade had knock-on cost, but it was never evidence for problem 2 — the assertion compares a hand-maintained module list against an AST scan, not label uniqueness
#4911 — open CI failure whose second of two saved signatures is exactly this ImportError under fix(kda): fall back to Cake when CuTe DSL predates cutlass.experimental #4667's version gate in environments with an older CuTe DSL; the first is an unguarded ModuleNotFoundError: No module named 'cutlass.experimental'. Not caused by "auto" routing (that fallback is the mitigation), but the same theme: "auto" and the strict backends have materially different totality
What #4483 got right, what it got wrong, and what is new since
#4483 raised this problem on 2026-08-12 and was closed on 2026-09-03 in favour of this issue. What changed in the three weeks between:
Resolved. Discoverability — #4483's fourth problem — is fixed: backend= landed on the top-level facade in #4605 and is forwarded to decode, so flashinfer.recurrent_kda(..., backend="cake") now reaches the strict Cake path. Two pieces of #4483's guidance were also followed: #4417 merged internal-only, and #4378 was closed in favour of #4445's public packed_kda_decode noun facade.
Superseded by events.#4483 described a single "opposite backend policy" between the two facades. That framing no longer fits. The facades now share one run_recurrent_kda and differ in their defaults, and the top-level default is not a safe superset — "auto" decode commits to Cake and raises rather than falling back, which is #4935. Its prefill description has also inverted, but by events rather than by error: #4483 described "auto" prefill as preferring frozen FlashKDA, which was correct on 2026-08-12 and is still what released v0.6.18 does (kda.py:260 hardcodes backend="cake"). CuTe-DSL prefill only arrived with #4605 on 08-20. Today "auto" prefers CuTe DSL with Cake as the fallback, and that fallback is nvidia-cutlass-dsl version-dependent (#4667).
New since #4483. The backend= parameter itself; the "auto" literal and the defect behind it; a training/backward family and a plan-and-run wrapper class (six top-level symbols, all decorated-but-untraced or classes, plus a private _kda_training_impl split); an SM120a prefill tranche that added three kda_kernels exports and a new subpackage; two more docs pages; and concrete fragmentation across seven open PRs, with #4709 as a counterexample that keeps both facades in step. Critically, none of that has shipped — see Release status — which is what makes the plan above cheap now and expensive after 0.7.
Unchanged and still reproducing. The duplicate recurrent_kda name and shared trace label; the top-level facade bypassing the decode wrapper; the stage:decode tag on a template that dispatches prefill; missing backend / seq_order / prefill_workspace trace fields; kda_decode.rst autosummarising recurrent_kda; the kda_kernels noun aliases; decode-only fi_trace coverage; and the prefill/decode test import split.
TL;DR. Two public functions named
recurrent_kdadisagree on their default backend, and the top-level default is not a safe superset of the other. The unification itself is not proposed for0.7— it touches released API and belongs on a normal deprecation schedule. Four things do want attention before the cut:backend="auto"decode raises instead of falling back, contradicting its own docstring, which says"auto"selects Cake for its native contract and "otherwise preserves CuTe DSL" (recurrent_kda.py:1873-1878). Fixing it is a restoration of the documented contract, not a behaviour change. The only release-critical engineering item here; fix proposed in #5037.0.7should not advertise them as stable API. The plan is to mark the unreleased surface experimental, which says plainly that these signatures may change as the unification below lands. See Holding action.backendliteral to one facade only. The cheapest intervention on this list, and it is open now.main.The rest is background and a proposed sequence.
Context
Status as of 2026-09-03 on FlashInfer
main(7a3c04f0).This supersedes #4483 (
[KDA] Unify flashinfer.recurrent_kda API, 2026-08-12). That audit is still worth reading as the historical record, but enough has changed in three weeks — abackend=parameter landed, the top-level surface grew from 3 symbols to 10, and a release was cut in the middle of it — that patching it in place would have been more confusing than restating it. Relationship to #4483 summarises what moved.FlashInfer has two decorated functions named
recurrent_kdaregistering the same trace template. Since #4483,backend=has been added to the top-level facade (#4605), which resolves the original discoverability gap. But the two facades now ship different defaults, and the top-level default is not a safe superset of the other:backend="auto"decode commits to a frozen Cake kernel and raises instead of falling back to CuTe DSL. That defect is filed separately as #4935; this issue covers the public Python API contract only.The KDA public surface also grew from the 3 top-level symbols #4483 audited to 10 — of which
v0.6.18shipped 4 — adding a training/backward family and a plan-and-run wrapper class that #4483 did not cover.Six of those ten symbols, and the entire
backend=parameter on the top-level facade, have never been in a release — see Release status. That is the most important input to the plan below: the parts of this design that are still wrong can mostly be fixed outright, but only until0.7is cut. Since that cut is close, Holding action proposes narrowing the advertised surface back to whatv0.6.18shipped rather than rushing the consolidation.This is a design / tracking issue. Related CAKE progress is tracked in #4254. The same problem class on the sibling GDN op is already tracked in #2687, with an in-flight implementation in #2706 — that pair is the closest precedent for the plan below and should be aligned with rather than duplicated.
Suggested implementation sequence
Steps are ordered by dependency, but the release window is the scheduling constraint: anything touching unreleased surface is cheap now and locked after
0.7.0.7?backend=toflashinfer.kda.recurrent_kda; wire Cake decode and prefill eligibility through itbackendliteral cover both facades or go in as an experimental backend under #4880"cudnn"tokda.py's enum and leaveskda_decode.pyuntouched — the same shape as #4779's"vibecuda". Raised on #4968 and #4779; #4779 also AOT-registers its backend, which #4880's policy forbids for an experimental onemain)"auto"as the top-level default, keeping the decode facade's released"cute-dsl"default untouched"auto"decode is total — see remedy 2 above"auto"decode kernel swap againstv0.6.18's effective"cute-dsl"routing on the intended serving path"auto"in step 2. Measured: Cake wins at every batch size under CUDA-graph capture; the eager small-batch loss is host dispatch_kda_training_implsplitis_experimentalkeeps them out of the stable inventory, so 3b-i removes the urgencyRecurrentKDAPrefillWrapperper-method onplan/run, plus the five training/backward symbols only if #4965 has not landedtests/experimental/, leaving the three CuTe-DSL-numerics tests that merely drive the wrapper in the stable lanekda_backward.py+_kda_training_impl.pybehind thin gated entry points, move their tests totests/experimental/, de-register the tenaot.pysites, file the owner / tracking issue / graduation plan0.7__all__half is done in #5040 (open), which drops the three unreleasedkda_kernels.__all__exports and restores__all__to byte-identical withv0.6.18; a zero-consumer change, since dispatch reads them as attributes and every doc, test and benchmark uses the deeperkda_kernels.sm120_prefillpath. Still open: the explicitbackend=literal and gating the automatic branch withexperimental_auto_backends_allowed()__all__half; the gate can follow"auto"prefill routing on CC 12.0 and mostly by raising (see remedy 2 above), so it wants its own PR validated on a CC 12.0 device_dispatch_recurrent_kda_decodeused by both facades; decide which facade owns therecurrent_kdatracekda_decode.recurrent_kdaat that dispatcher with omitted→"cute-dsl"+ deprecaterecurrent_kdatraces distinguishable labels; extend the trace template and retag stagetest_recurrent_kda.py, Cake export tests, and fi_trace to the canonical facade; add top-levelfi_apitrace coveragekda_kernelsnoun aliases (un-shadows submodules); consolidate docs pages; document or internalizeFLASHINFER_KDA_T1_FAST_PATHThe time-critical items are 0b, 1, 2, 2b, 3b-i and 3c; all are confined to
main-only surface. Everything else concerns released API and can proceed on a normal deprecation schedule.Stated plainly: if step 3b-i lands, the release-critical engineering reduces to two items — step 1, and step 3c's disposal of the three unreleased
kda_kernelsSM120a exports, because the top-levelbackend=parameter and its default would then be the sole unreleased inference surface that a0.7cut freezes — and step 2 cannot be settled without it, since nobackendvalue is safe across both phases until"auto"decode is total. Step 0b is one comment; step 3's subject disappears with 3b-i.The benchmark in step 2b is not a formality. On a B200 (
H=32,K=V=128,T=1, eager, 200 iterations)"auto"decode selects Cake and crosses over against"cute-dsl"at aboutB = 35: below it"auto"is up to 3.1x slower, above it up to 7.5x faster (0.32x atB=4, 1.90x atB=64, 7.53x atB=256). Cake's flat ~108 µs is host dispatch rather than device time, and under CUDA-graph capture the picture inverts completely: Cake wins at every batch size, by 3.7x to 7.9x. So the eager small-batch loss is ~75 µs of Python-side enqueue that capture erases, and no size condition on the gate is warranted — the default is right for the graph-captured serving path it exists for. What remains worth confirming is that the intended consumers do capture.If step 1 slips, the contingency is not to ship
"auto"as-is. Disabling the speculative"auto"→ Cake decode gate is a one-line change toauto_unbounded_softplus_candidate(recurrent_kda.py:2063-2064) that makes"auto"decode identical to"cute-dsl"— bit-exact on every trigger class with a non-zero pool and non-identity indices,backend="cake"strictness intact, and exactly the routing a defaultedv0.6.18call took. It costs the Cake decode win at serving batch sizes for one release and requires retargeting six parametrizations oftest_t1_unbounded_softplus_auto_route_*tobackend="cake". Land step 1; keep this in reserve.Current public surface (main)
Full inventory: both facades, the ten top-level symbols, and the docs split
@flashinfer_apiflashinfer.recurrent_kdaflashinfer/kda.pyrecurrent_kda_trace)backend(default"auto"),seq_order,prefill_workspace,state_checkpoints,checkpoint_cu_starts,checkpoint_every_n_tokens_is_plain_multi_token_prefill). Plain multi-token prefill: under"auto"tries SM120a CuTe → SM100-family CuTe DSL BT=16 → frozen Cake fallback. Otherwise decode/spec →run_recurrent_kdavia raw_run_recurrent_kda(still bypasses decoratedkda_decode.recurrent_kda), forwardingbackend=flashinfer.kda_decode.recurrent_kdaflashinfer/kda_decode.pybackend(default"cute-dsl")run_recurrent_kda(backend=…). Default"cute-dsl"→ CuTe DSL, never Cakeflashinfer.packed_kda_decodeflashinfer/kda_decode.pypacked_kda_decode_trace)T=1,H=12,K=V=128; strict Cake (#4445)flashinfer.fused_kda_decodeflashinfer/kda_decode.pyfused_kda_decode_trace)flashinfer.RecurrentKDAPrefillWrapperflashinfer/kda.pyplan()/run()runforcesbackend="cute-dsl"flashinfer.RecurrentKDAPrefillWorkspaceflashinfer/kda_prefill.pyflashinfer.recurrent_kda_backward,RecurrentKDABackwardWorkspaceflashinfer/kda_backward.pyflashinfer.recurrent_kda_training_forward/_backward,RecurrentKDATrainingContextflashinfer/kda_training.py(re-export of private_kda_training_impl.py)Top-level
__init__.pyre-exports 10 KDA symbols (__init__.py:165-184):recurrent_kdaandRecurrentKDAPrefillWrapperfromkda;fused_kda_decodeandpacked_kda_decodefromkda_decode;RecurrentKDAPrefillWorkspacefromkda_prefill;recurrent_kda_backwardandRecurrentKDABackwardWorkspacefromkda_backward; andrecurrent_kda_training_forward,recurrent_kda_training_backward,RecurrentKDATrainingContextfromkda_training.flashinfer.kdais reachable as an attribute (the package-rootfrom .kda import …binds the submodule) but is not an intentional re-export.flashinfer.kda_kernelsstill re-exports three noun aliases (recurrent_kda,fused_kda_decode, and nowpacked_kda_decode) alongside therun_*backends. Its__all__has grown from 6 entries to 9 sincev0.6.18; all three additions are the SM120a prefill helpers, which are unreleased. Two of the three are not merely redundant — they shadow a submodule of the same name, so attribute access cannot reach the module:In-repo consumers that need the shadowed modules work around this with
importlib.import_module. That makes the alias cleanup a correctness fix for the first two, not cosmetics. Separately,kda_kernels/__init__.pydoes now carry an explicit note thatpacked_kda_decode_cuteis an internal implementation module — the desired precedent, applied inconsistently with the shadowing noun aliases directly above it.Docs are split across five pages:
docs/api/kda.rst,kda_decode.rst,kda_prefill.rst, and the newkda_backward.rstandkda_training.rst(both deleted by #4965).kda_decode.rststill autosummarizesrecurrent_kda,packed_kda_decode, andfused_kda_decode.kda_prefill.rstdocumentsRecurrentKDAPrefillWrapper, which is defined inkda.py.In-repo usage today: Prefill tests/benchmarks use the top-level /
kdafacade (bench_recurrent_kda_prefill.py,tests/kda/test_recurrent_kda_prefill.py). Decode paths — including Cake export tests and the main CuTe DSL correctness suitetests/kda/test_recurrent_kda.py, plusbench_recurrent_kda.py,bench_recurrent_kda_decode_export.py, andbench_recurrent_kda_decode_split_sweep.py— importflashinfer.kda_decode.recurrent_kda. The newertests/kda/test_packed_kda_decode_cute.pyuses top-levelflashinfer.recurrent_kda, so the codebase now exercises both facades for decode-shaped calls. fi_trace fixtures still pin only the decode facade.The split is also why CI does not catch #4935, though the mechanism is subtler than "the top-level default is untested". Across
tests/kda/, the"auto"gate matches on 6 of 190run_recurrent_kdacalls, all six with an explicitbackend="auto"and all from one test that asserts Cake is selected — so it cannot detect a missing fallback. The three defaulted top-level calls intest_packed_kda_decode_cute.pyall miss the gate, two onlower_bound=-5.0and the third onuse_gate_in_kernel. The gap is a missing negative test, not missing coverage, which makes it cheap to close.Release status: what is actually shipped
Most of the surface this issue describes has never been in a release. That widens the design freedom considerably, and it also puts a clock on it.
v0.6.18(published 2026-08-29, tag commit69ff11fc) is the current stable release. It is not an ancestor ofmain: the release branch was cut frommainat61a6c651on 2026-08-19 (author date01:03Z; 08-18 in US-Pacific, which is how it appears ingit loglocally), then took cherry-picks through 08-28. So "did it ship?" cannot be inferred from merge dates — the classification below comes from reading thev0.6.18tree directly.v0.6.18recurrent_kda(phase-neutral facade, nobackend=)fused_kda_decodepacked_kda_decodeRecurrentKDAPrefillWorkspaceRecurrentKDAPrefillWrapperrecurrent_kda_backward,RecurrentKDABackwardWorkspacerecurrent_kda_training_forward/_backward,RecurrentKDATrainingContextSo
v0.6.18exports 4 KDA symbols wheremainexports 10.No other channel carries the newer surface either. PyPI serves
flashinfer-python0.6.18 as latest (plus a0.6.18rc10prerelease, checked and clean); therelease-v0.6.18branch head is the tag itself; there are no0.6.18.post*releases, which is informative rather than accidental sincev0.6.16had four. Nightly tagging stops atnightly-v0.6.18-20260819— one day before #4605 — and nightlies trackmainrather than the release branch, which is what makes that check meaningful. No nightly contains a top-levelbackend=or theauto_unbounded_softplus_candidategate.Per-item breakdown: what is changeable outright vs. what carries a compatibility constraint
Unreleased — changeable outright, no deprecation cycle
backend=parameter on the top-level facade, in its entirety. Releasedflashinfer.recurrent_kdahas no such parameter; the onlybackend=in the releasedkda.pyis an internal hardcodedbackend="cake"prefill call."auto"as the top-level default is therefore unreleased."auto"as a literal anywhere. Releasedkda_decode.recurrent_kdaandrun_recurrent_kdaboth declareLiteral["cute-dsl", "cake"] = "cute-dsl", and theauto_unbounded_softplus_candidategate is absent from the released kernel module. [Bug]recurrent_kdabackend="auto"decode commits to Cake and raises instead of falling back to CuTe DSL #4935 has never shipped, so it can be fixed by redesign rather than by a compat-preserving patch.RecurrentKDAPrefillWrapperand the entire training/backward family (five symbols), with their moduleskda_backward.py,kda_training.py,_kda_training_impl.py,_kda_training_dispatch.py,kda_prefill_cute.pyand docs pageskda_backward.rst,kda_training.rst. The bare-decorator/no-trace choice and the private_*_implsplit (problem 8) can simply be corrected. The unreleased footprint is wider than the public symbols suggest: 15 KDA-related Python modules exist onmainand not inv0.6.18, the rest beingkda_kernels/kda_chunked_bt16.py, the fourkda_kernels/sm120_prefill/files, and fivejit/modules (_kda_jit_common.py,cake_kda.py,cake_kda_decode.py,flash_kda_backward.py,flash_kda_training.py).21a95d73, but only its sharedis_cute_dsl_experimental_availableprobe (14 lines incute_dsl/utils.py) — the KDA routing did not ship, and the release engineer said so in the commit message: "0.6.18 has no kda_prefill_cute / feat(kda): add CuTe DSL recurrent prefill backend #4605 auto CuTe prefill path (public recurrent_kda already uses Cake), so only the shared is_cute_dsl_experimental_available probe is cherry-picked." That is upstream confirming this section's thesis in the release branch's own history.kda_kernels.__all__grew from 6 entries to 9, and all three additions are SM120a:can_implement_kda_prefill_sm120,clear_kda_prefill_sm120_caches,run_kda_prefill_sm120. Behind them sits a newflashinfer/kda_kernels/sm120_prefill/subpackage (decomp.py,fused.py,runtime.py), the_sm120_kda_prefill_*eligibility and rejection machinery inkda_prefill.py,RecurrentKDAPrefillWorkspace._sm120_state, and the"auto"prefill route that tries SM120a first (kda.py:302-306). None of it is inv0.6.18.Released — real compatibility constraints
Note that the core complaint of this issue is in the released set, so consolidation still needs a deprecation story:
recurrent_kdafacades shipped under the same name, and both were already decorated withrecurrent_kda_trace— the duplicate name and the trace-registry collision (problems 1 and 2) are released._is_plain_multi_token_prefilland the top-level facade's use of it both shipped, while the decode facade never classified phase.kda_decode.recurrent_kda's keyword-onlybackend: Literal["cute-dsl", "cake"] = "cute-dsl".mainhas since widened that enum to include"auto"(feat(cake_kda): add native unbounded-softplus Kimi-Linear kernels #4535), which is purely additive; changing the default would be a breaking behavior change.seq_order,prefill_workspace, and the checkpoint kwargs on the top-level facade.kda_kernelsnoun aliases and their submodule shadowing (shim item 4).FLASHINFER_KDA_T1_FAST_PATH(problem 6).Problem 3(a) — the default divergence — is the one part of problem 3 that is unreleased, precisely because the released top-level facade had no
backendat all.Timing
mainis still stamped0.6.18, and no0.7branch or rc tag exists yet. Recent stable cadence is roughly two to three weeks (v0.6.172026-08-11,v0.6.182026-08-29), so the next cut is plausibly mid-September. Once0.7shipsbackend="auto"as the top-level default, changing it presumably owes users a deprecation cycle rather than being a straight correction, and the cheap fixes here turn into expensive ones. The unreleased-surface decisions should land before the next release cut.Holding action: what
0.7should advertiseThe plan is not to advertise the new KDA surface as stable in
0.7. Six of the ten top-level symbols have never shipped, and several are ones this issue proposes to change. Marking them experimental states that directly: they remain usable, but their signatures may change as the unification below lands, and nobody acquires a compatibility claim on them in the meantime.Update (2026-09-04): five of those six are now proposed for removal rather than marking. Per @YangXu1990uiuc's comment, FlashInfer and cuDNN Frontend are splitting KDA by phase — inference here, training there, where
cudnn.linear_attention.ops.kimi_delta_attentionalready provides fwd+bwd with autograd and the CAKE training kernels will land as an engine under it. #4965 implements the removal ofrecurrent_kda_training_forward,recurrent_kda_training_backward,RecurrentKDATrainingContext,recurrent_kda_backwardandRecurrentKDABackwardWorkspace. Since none has shipped and a GitHub-wide search finds no external caller, removal before the cut needs no deprecation cycle, and it disposes of problem 8 and theaot.pyregistrations outright instead of relocating them. That leavesRecurrentKDAPrefillWrapperas the only unreleased top-level symbol needing the experimental treatment. #4965 is still a draft pending its cuDNN Frontend counterpart (NVIDIA/cudnn-frontend#912), so the marking below remains the fallback for the training family if the cut lands first. The rest of this section is written for that fallback; ordinary prefill, including them128_n16_checkpointvariant and the checkpoint kwargs that serve SGLang prefix caching, is untouched either way.The mechanism is #4880, which merged on 2026-09-04. Marking a symbol is non-breaking:
@flashinfer_experimental_apiemits anExperimentalWarningonce rather than raising, and its docs are explicit that calling an experimental API is itself the opt-in, requiring no environment variable. What the env varFLASHINFER_ALLOW_EXPERIMENTAL_AUTO_BACKENDS=1gates is automatic selection — whetherbackend="auto"may pick a backend marked@experimental_backend.traceis optional for experimental APIs andis_experimentalexcludes them from the stable trace inventory, which incidentally sanctions the training family's bare-decorator state (problem 8) rather than leaving it an inconsistency.Applying it means following #4880's own guidelines rather than bolting a decorator onto the current layout. Concretely, for the KDA surface:
Decorate methods, not classes.
@flashinfer_experimental_apireturns a function, so applying it to a class silently replaces the class with a function and breaksisinstance, subclassing and Sphinxautoclass— anddocs/api/kda_prefill.rstautosummarizesRecurrentKDAPrefillWrapper, so that would be a visible docs regression. Per-method decoration avoids this and is already the codebase's convention for class-based public ops: 98 indented@flashinfer_apisites, includingsparse.py:1224whereBlockSparseAttentionWrapper.runcarries a full trace template, against zero applications to a class. That covers all six symbols:RecurrentKDAPrefillWrapperplan/runrecurrent_kda_backwardrecurrent_kda_training_forwardrecurrent_kda_training_backwardRecurrentKDABackwardWorkspacerecurrent_kda_backwardRecurrentKDATrainingContextWorth adding an
isinstance(f, type)guard to the decorator so the degradation is an error rather than silent; raised on #4880. Note that per-method decoration warns but does not document:RecurrentKDABackwardWorkspaceandRecurrentKDATrainingContexthave their ownautosummaryentries and would render as ordinary stable pages, so each needs an explicit.. warning::in its class docstring.Marking and containment are separate pieces of work, and only marking belongs before the cut. Decorating the symbols is hours of work with no kernel, wheel or CI-lane change. Containment is not: #4880's
CONTRIBUTING.mdrequires experimental tests to live intests/experimental/, and itspytest.inisetsnorecursedirs = test_helpers experimental, so complying movestests/kda/test_recurrent_kda_backward.pyandtest_recurrent_kda_training.py— 1,522 lines of Blackwell coverage — out of the default lanes days before a cut. De-registering also breakstest_aot_registers_flash_kda_backward, which pins the spec list exactly (tests/jit/test_flash_kda_backward_jit.py:206-213), and the real footprint is tenaot.pysites plus three tests, not four entries: imports at:98-99, capability probes at:553-557, the four appends at:674-679, anddetect_sm_capabilitiesat:1235-1242. The relocation is also not a file move — the containment rule wants a thin core entry point, andflashinfer/kda_training.pyis currently a bare re-export, so the 15-parameter signature has to be restated in core, wheretests/kda/test_recurrent_kda_training.py:86-121pins it verbatim and a*args, **kwargshandoff fails. First-call JIT for training users is roughly 91 s per cache directory, measured. None of that is an argument against containment — it is the agreed destination — but nothing in the0.7contract depends on file placement, so it should land in the first release after.RecurrentKDAPrefillWrapperis exempt either way: it is host-side planning that ends in a call to the stablerecurrent_kdafacade, so it has no backend implementation to contain.Relocate and de-register, per the containment rule — the fallback path only, superseded if #4965 lands. #4880's README is explicit that a core entry point may include only the signature, shared validation, the gate check, backend selection and a handoff, and that "if the API is experimental, its implementation lives here too, regardless of the implementation's maturity — this keeps removal a one-directory deletion". So
kda_backward.pyand_kda_training_impl.py(~3,000 lines) move underflashinfer/experimental/behind thin gated entry points, and the tenaot.pysites enumerated above come out, leaving the training and backward kernels JIT-only as the policy requires. Two consequences to expect rather than to debate: wheel contents change, and training users take first-call JIT latency. Both are the normal price of the experimental track, and both reverse cleanly when these APIs graduate.One genuine convergence, worth aligning deliberately.
FLASHINFER_ALLOW_EXPERIMENTAL_AUTO_BACKENDSexists precisely so thatbackend="auto"declines to select experimental backends — which is exactly the shape of the SM120a prefill question. The unreleased surface splits by architecture, and the two halves want different halves of the policy. The six top-level symbols are Blackwell-datacentre only:recurrent_kda_backwardand the training trio gate on_SUPPORTED_COMPUTE_CAPABILITIES = ((10, 0), (10, 3))(kda_backward.py:35,_kda_training_impl.py:22), andRecurrentKDAPrefillWrapperis documented "Compute capability 10.0 and 10.3 only". Those are API surface. The SM120a prefill tranche is backend surface reached through a stable API — the shape@experimental_backendis built for. It cannot use that decorator yet:@experimental_backendmarks a@backend_requirementchecker, and KDA has none, so it only becomes available after the migration in remedy 2. Until then the hand-rolled equivalent iswarn_experimental_backend_onceon an explicit branch plusexperimental_auto_backends_allowed()guarding the automatic one atkda.py:302— notrequire_experimental_auto_backends, which raises rather than falling through. Gating SM120a also changes"auto"prefill routing on CC 12.0, and mostly by raising. Of three prefill spellings that work on CC 10.0, only the packed one still succeeds once SM120a is gated off: a fixed-Tcall (q=[B, T>1, H, 128], nocu_seqlens) falls through to the generic path and then raisesDecode only supports T=1 without cu_seqlens, and a call passingprefill_workspace=raisesValueErrorbecause SM120a is the only CC-12.0 backend that accepts a workspace (prefill_workspaceatkda_prefill.py:3627, the raise atkda.py:473). Neither is a regression againstv0.6.18, which had no CC-12.0 prefill backend at all, but the second contradicts a "falls through rather than raises" reading. Worth noting #4880's stated default here is to raise with a message naming the variable, so a fall-through is a deliberate divergence. There is no SM121 question: the gate is exactly_SM120_KDA_SUPPORTED_COMPUTE_CAPABILITIES = {(12, 0)}(kda_prefill.py:3594, mirrored bykda.py:325), so CC 12.1 is served by no KDA prefill backend, released or not.The architecture predicate is enumerated in eight Python places, four of which independently spell the same pair.
((10, 0), (10, 3))appears verbatim atkda_backward.py:35,_kda_training_impl.py:22,kda_prefill_cute.py:34andkda_prefill.py:43— the last two being the prefill entry-point gates. Alongside them:_FLASH_KDA_DECODE_ARCH_BY_COMPUTE_CAPABILITY(recurrent_kda.py:70-75) for decode, and(12, 0)three times over for SM120a —_SM120_KDA_SUPPORTED_COMPUTE_CAPABILITIES(kda_prefill.py:3594),SM120_CAPABILITY(kda_kernels/sm120_prefill/runtime.py:138) andDEVICE_CC(sm120_prefill/fused.py:99). #4710 is a fair illustration rather than a counterexample: it widened the fourcsrc/kda/*.cuhruntime guards to admit CC 10.7 and refreshed the JIT cache keys that edit invalidates, and its scope note states plainly that the Python entry points still restrict KDA to 10.0/10.3 so end-to-end Rubin support is separate work. The observation is only that the separate work is spread across five sites, and that while"auto"raises on an unmapped capability (#4935) they have to move in lockstep with the kernels.Marking does not fix #4935.
"auto"is the top-level default and the defect is in its decode routing, so that needs the change in step 1 regardless.The marking and the bug fix are independent, and only the bug fix is release-critical. Marking is worth doing before the cut because it is what keeps the unification below cheap, but nothing breaks if it slips. #4935 is different: it is the behaviour of the default path, so shipping it means shipping a broken default.
How we got here
Timeline of the merged PRs that produced the current shape
flashinfer.kda.recurrent_kdaand switched the top-level export to it; automatic frozen prefill dispatchbackend="cake"tokda_decode.recurrent_kdaand therun_recurrent_kdabackend — updating the older decode facade, not the new top-level onefused_kda_decode(reasonable separate op: conv + recurrence + norm)packed_kda_decodeas a public noun-form facade (trace + top-level export + docs), superseding #4378. Also added checkpoint kwargs to the top-level facadebackend="auto"|"cute-dsl"|"cake"toflashinfer.kda.recurrent_kdaand wired CuTe DSL recurrent prefill. This is Step 1 of #4483's proposed sequence"auto"/"cute-dsl""auto"fall back to Cake when CuTe DSL predatescutlass.experimental(<4.7), and made explicit"cute-dsl"raiseImportErrorthere — introduced version-dependent routing. The strict path is currently red in CI (#4911)"auto"tokda_decode.recurrent_kdaandrun_recurrent_kda, selecting Cake for the equal-head/D128/T=1unbounded-softplus contract. Source of #4935The
v0.6.18branch cut falls between #4445 (08-16, released) and #4605 (08-20, unreleased). Everything from #4605 onward in that table — the top-levelbackend=parameter, both CuTe DSL prefill backends, the version gate, and"auto"decode — ismain-only.Concrete problems
Duplicate canonical name. Two
@flashinfer_apiwrappers share the namerecurrent_kdaand the same trace template, and are not the same object (flashinfer.kda.recurrent_kda is flashinfer.kda_decode.recurrent_kda→False). The top-level facade still bypasses the decode facade by calling raw_run_recurrent_kda, so decode through the top-level path is never seen by the decode wrapper's own logging/tracing. (It does now forwardbackend=, so the original "never setsbackend" finding no longer applies.)Trace registry collision. Both facades register in
_TRACE_REGISTRYunder the label'recurrent_kda'. The count depends on import scope — plainimport flashinfershows two colliding labels, and importing every submodule shows four:(Reproducing the second line needs a sweep that tolerates import failures; a naive import-everything loop raises before finishing.)
These are not all the same defect, and the distinction matters for the fix:
recurrent_kdakda,kda_decode) — the duplicate-facade problem this issue is aboutgdn_decodecute_dsl_fused_moe…_fused_moe,…_fused_moe_nvfp4) in one module sharing a templateblock_sparse_runOnly
recurrent_kdais two same-named public facades; the other three are deliberate template reuse that merely produces a non-unique label. Any registry-uniqueness fix should distinguish these cases rather than treat them alike.No test catches this. The registry's duplicate key is
(func.__module__, func.__qualname__, label)(tests/trace/template_registry.py:100-103), and the two facades differ in__module__, so the shared label never trips the guard. The collision is not a failing test — it is unpoliced, which is the weaker-looking but more durable argument for fixing it.recurrent_kda_tracehas since gainedssm_state_indices,state_checkpoints,checkpoint_cu_starts, andcheckpoint_every_n_tokens, but still does not modelbackend,seq_order, orprefill_workspace, and is still taggedstage:decodedespite dispatching prefill. fi_trace fixtures still pinfi_api:flashinfer.kda_decode.recurrent_kdaonly, so top-levelfi_apistrings remain untested.The two facades disagree on the default, and the top-level default is not a superset. Both delegate decode to the same
run_recurrent_kda, so for a fixedbackendvalue the decode kwargs are identical. Two things still diverge:(a) The default value. Verified by intercepting the forwarded value:
backenddefaultT=1decodeflashinfer.recurrent_kda"auto"ValueErrorif Cake cannot serve it ([Bug]recurrent_kdabackend="auto"decode commits to Cake and raises instead of falling back to CuTe DSL #4935)flashinfer.kda_decode.recurrent_kda"cute-dsl"Because
"auto"commits to Cake without falling back (#4935), the top-level default is not a superset of the decode facade's — and the two accepted sets are in fact incomparable, since one guard is skipped under"auto"that"cute-dsl"enforces. The same call can succeed through one import path and raise through the other, in both directions. This is the strongest argument for consolidation, but it is a defect to fix in [Bug]recurrent_kdabackend="auto"decode commits to Cake and raises instead of falling back to CuTe DSL #4935, not something to design around here.(b) Phase classification runs before delegation.
flashinfer.kda.recurrent_kdacalls_is_plain_multi_token_prefillfirst, which classifies on shape only (kda_prefill.py:315-327): nonum_spec_tokens, and either fixedT > 1or packedtotal_tokens > num_sequences. There is no head-dim, dtype, or architecture guard. So a packed multi-token decode call is claimed by the prefill branch:The decode facade never runs this classifier. The
"auto"outcome here is gate-mode dependent — with the raw Kimi-Linear softplus gate (use_gate_in_kernel=True,lower_bound=None) the same call raises instead, because it then matches the"auto"Cake gate discussed in [Bug]recurrent_kdabackend="auto"decode commits to Cake and raises instead of falling back to CuTe DSL #4935. Either way the conclusion holds and in fact strengthens: no singlebackendvalue reproduces the decode facade's current behavior. This is the constraint that shapes the shim design below.Discoverability.Resolved by feat(kda): add CuTe DSL recurrent prefill backend #4605.flashinfer.recurrent_kda(..., backend="cake")now exists and reaches the strict Cake decode path.Docs / traces lag the facade. Five Sphinx pages now describe pieces of one recurrent API;
kda_decode.rststill autosummarizesrecurrent_kda;kda_prefill.rstdocuments a class defined inkda.py. Both thekda_decode.pydocstring anddocs/api/kda_decode.rstdescribe an"auto"→ CuTe-DSL decode fallback that does not exist (see [Bug]recurrent_kdabackend="auto"decode commits to Cake and raises instead of falling back to CuTe DSL #4935).kda_decode.rstdocuments the Cake export set thoroughly — the exported bodies, gate modes, per-Tvalue-row splits and per-architecture split policies — but not the layout boundary: whether a given call is accepted depends on a 259-line selector over strides, pointer alignment and int32-overflow products that no user can predict from the prose. Once"auto"falls back correctly this is a performance-predictability problem rather than a correctness one, but it should be stated somewhere. The trace template gaps andstage:decodetag are covered in problem 2.An environment variable participates in public dispatch, documented only for contributors.
FLASHINFER_KDA_T1_FAST_PATH(default"1",recurrent_kda.py:139) selects a differentT=1decode route. It appears inCLAUDE.md:607but nowhere indocs/api/, so users of the public API cannot discover it. Which kernel a given publicbackend=value resolves to therefore depends on process environment. Any unified contract needs to either document this as public API or make it internal-only.Fragmentation is now happening, not hypothetical. Open PRs are each adding surface in a different style, concurrently editing the same two facades:
flashinfer/kda_vibecuda.pyand widens the top-level enum toLiteral["auto","cute-dsl","cake","vibecuda"]onkda.pyonly — a fourth backend literal that the decode facade will not accept, pushing the two signatures further apartkda.pyandkda_prefill.pyand adds ajit/flash_kda_nvrtc.pycompilation pathfused_kda_decode_packed— a fourth*_kda_decodenounkda.py/kda_prefill.pyprefill backendskda_decode.pyfeat(vibecuda_kda): add SM100/SM103 recurrent KDA prefill backend #4779 is the immediate instance: it adds
"vibecuda"tokda.py's enum only, which is the divergence this issue is about — no comment on the kernel work, just the one line that would need to cover both facades.Worth citing as the counterexample: #4709 adds
disable_state_updatewith backend-specific raises but updates bothkda.pyandkda_decode.py, keeping the two signatures in step. That is the behavior this issue wants to make the default expectation.New surface outside the original scope.
kda_backward.pyandkda_training.pyintroduced a paired-function training family, andRecurrentKDAPrefillWrapperintroduced a plan-and-run class convention. Neither is covered by the "noun-form facade vs internalrun_*" rule below, so that rule needs cases for stateful plan/run wrappers and paired forward/backward ops. Two further inconsistencies came along with them:@flashinfer_apiwith no trace template, unlike everyrecurrent_kda/*_kda_decodefacade. So "public op" currently means "decorated" in two incompatible ways, and the training family is invisible to fi_trace.flashinfer/kda_training.pyis a 23-line re-export of privateflashinfer/_kda_training_impl.py— a module-layout convention used nowhere else in KDA (kda_backward.pykeeps its implementation inline). There is a stated reason: the module also defines_get_training_module, an "internal indirection seam used by the training API tests". That justifies the seam; whether it also justifies the split is worth a sentence from the author. Either way it is cosmetic next to the missing trace templates alongside it.Adjacent: #4896 (production KDA prefill for SGLang prefix caching).
Non-goals
fused_kda_decodeorpacked_kda_decodeintorecurrent_kda— different serving contracts (packed QKV, conv, RMSNorm; different return types).recurrent_kdabackend="auto"decode commits to Cake and raises instead of falling back to CuTe DSL #4935 here. It is a behavior defect with its own regression test; this issue depends on it but should not absorb it.Proposed target (backwards-compatible)
Canonical API
One public recurrent inference entrypoint — this now exists and already accepts
backend=:Current
backendsemantics, as actually implemented onmain:backend"auto"(default)nvidia-cutlass-dsl<4.7). Decode: commits to Cake when the equal-head/D128/T=1unbounded-softplus gate matches, raising rather than falling back if Cake cannot serve it (#4935)"cute-dsl"ValueErroron an unsupported plain-prefill contract,ImportErrorwhen the DSL predates 4.7"cake"Target semantics:
"auto"becomes a true superset — never raising where"cute-dsl"or"cake"would have succeeded. The plan below assumes that fix has landed.Because none of this parameter has shipped, the fix should be the structural one described in #4935 — resolve backend eligibility before state normalization, so one resolved decision drives both the state convention and the launch — rather than the cheaper "keep the speculative gate and hand-synchronise a second predicate" variant. There is no released behavior to preserve, so there is no reason to accept two predicates that must be kept in step by hand; that fragility is what produced the defect.
The remaining work is not adding this parameter — it is reconciling the second facade against it. Keep
seq_order/prefill_workspace/ checkpoint kwargs on this facade only (all three are released there).Compatibility shims
flashinfer.kda_decode.recurrent_kdacannot simply forward toflashinfer.kda.recurrent_kda. No singlebackendvalue preserves today's behavior, because the two facades differ in both the default and the phase classifier (problem 3):"cute-dsl""auto"T=1Kimi-Linear equal-head/D128recurrent_kdabackend="auto"decode commits to Cake and raises instead of falling back to CuTe DSL #4935)num_spec_tokens=NoneSo the shim must bypass phase classification and target a shared decode dispatcher directly, rather than re-entering the phase-neutral facade:
kda.recurrent_kdainto one internal_dispatch_recurrent_kda_decode(...), used by both facades, giving decode a single contract regardless of entry point. Note this does not by itself fix tracing: an internal dispatcher is undecorated by construction, so problem 1 additionally requires deciding which facade owns therecurrent_kdatrace (see the docs/traces section).kda_decode.recurrent_kdaforwards to that dispatcher with omittedbackend→"cute-dsl", preserving today's kernel selection exactly."cake"and"auto"unchanged.DeprecationWarningpointing atflashinfer.recurrent_kda(soft; keep for ≥1 minor release).seq_order/prefill_workspaceon the shim — those are prefill-only and the shim deliberately does not reach prefill.Decide the top-level default explicitly, before
0.7. This is a policy choice separate from [Bug]recurrent_kdabackend="auto"decode commits to Cake and raises instead of falling back to CuTe DSL #4935: even once"auto"stops raising, it selects a different decode kernel than"cute-dsl"would. Today, no single value is a safe default for a phase-neutral facade — and that is the sharpest evidence for this whole issue. The two phases have opposite totality properties:backend="auto"backend="cute-dsl"initial_statepool that"auto"acceptsImportErrorbelownvidia-cutlass-dsl4.7 (kda.py:366-372) — the second of the two failures in [Bug][v0.6.18] test_recurrent_kda_prefill.py ModuleNotFoundError: No module named 'cutlass.experimental', ImportError: backend='cute-dsl' requires nvidia-cutlass-dsl>=4.7.0 (cutlass.experimental); #4911Neither value is total on either phase, and on decode the two accepted sets are incomparable rather than nested (see #4935, Non-comparability);
"cake"is strictly narrower than both. The prefill raise says so in its own message: "backend='cute-dsl' requires nvidia-cutlass-dsl>=4.7.0 (cutlass.experimental); backend='auto' falls back to Cake." So defaulting the top-level facade to"cute-dsl"would break defaulted prefill calls on older CuTe DSL, and defaulting it to"auto"ships the known-broken decode default. A singlebackendknob on a phase-neutral facade is being asked to encode two independent phase-specific decisions, which is problem 3(b) restated as a concrete release blocker.That resolves the choice rather than blocking it:
"auto"is the right top-level default, but only once step 1 lands. It is already the only value with fallback semantics on both phases; step 1 makes those semantics true on decode as well as prefill. Without step 1 there is no safe default at all, which is the strongest reason to treat it as release-critical rather than merely desirable.The top-level default is unreleased, so it can be set either way with no deprecation — but note that in
v0.6.18the top-level facade had nobackendat all, so a defaulted top-level decode call reachedrun_recurrent_kda's own"cute-dsl"default. That call has already changed kernel between the release andmain— and on the inputs in #4935 it now raises where thev0.6.18spelling of the same call worked, which makes step 1 a regression fix against the last release rather than merely a correction to unreleased surface. Ratifying"auto"is a silent kernel swap for existing top-level callers, not a neutral blessing of the status quo, and it should be benchmarked rather than assumed.Adopt
@backend_requirementrather than hand-rolling the dispatch.utils.py:1185already provides the repo's standard mechanism: per-backend eligibility checkers plus an"auto"heuristic, resolved by computing the set of suitable backends first — filtered on compute capability and problem size — and raising only when that set is empty. It is applied at 44 sites across 15 modules (gemm/,grouped_mm/,fused_moe/,quantization/,deep_gemm.py) and at zero sites in KDA, and it composes with@flashinfer_apiand a trace template at 43 of those 44 (e.g.attn_scores.py:1020-1021). Its resolve-then-dispatch ordering is what makes commit-then-check hard to write by accident, which is why it is the right target shape rather than a bespoke selector, and it gives the phase asymmetry above a natural home as two decorated entry points with their own checker sets.Two limits to plan around, though, which is why this is post-
0.7and not a wholesale replacement. Checkers run before normalization (utils.py:1412-1459validates and then calls; there is no post-normalization hook), so they can host KDA's contract predicates — trigger classes 1-3 — but not its layout predicates, which are evaluated against re-viewed and synthesized tensors that do not exist until:2263-2337and which produce trigger class 4. The decorator also injects a publicskip_checkkwarg whoseTruepath bypasses all checking. So the realistic shape is contract predicates in checkers and a late fallback for layout, not one mechanism for both.The decode facade's
"cute-dsl"default is released and is decode-only, so no prefill concern touches it and it should not be changed at all. Its enum already accepts"auto"(kda_decode.py:79, added by feat(cake_kda): add native unbounded-softplus Kimi-Linear kernels #4535), so nothing needs adding there. The two facades converge on"auto"only after step 1 — and if step 1 slips, they should not be converged at this cut.flashinfer.fused_kda_decode/packed_kda_decode: unchanged.flashinfer.kda_kernels.run_*: document as backends only; remove the three noun aliases (recurrent_kda,fused_kda_decode,packed_kda_decode) fromkda_kernels.__all__. This un-shadows the same-named submodules and lets in-repo consumers drop theirimportlib.import_moduleworkarounds. Extend the existingpacked_kda_decode_cute"not public API" note into a module-wide statement. These aliases shipped inv0.6.18, so the removal needs the same soft-deprecation window as item 1 even though it is a correctness fix.Docs / traces
docs/api/kda.rst);kda_decode.rstfocuses onfused_kda_decode+packed_kda_decodeplus a compatibility note for the oldrecurrent_kdaimport;kda_prefill.rstremains workspace/contract detail; consider movingRecurrentKDAPrefillWrapperonto the page for its defining module (Sphinx resolves it correctly today via.. currentmodule:: flashinfer.kda, so this is page organisation rather than breakage);kda_backward.rst/kda_training.rststay as the training pair."auto"fallback description inkda_decode.pyanddocs/api/kda_decode.rst(tracked in [Bug]recurrent_kdabackend="auto"decode commits to Cake and raises instead of falling back to CuTe DSL #4935).FLASHINFER_KDA_T1_FAST_PATH.recurrent_kda_traceforbackend,seq_order, andprefill_workspace, retag the stage now that the template covers prefill, and give the two facades distinguishable registry labels (or collapse them once the shim lands). Treat thegdn_decode/cute_dsl_fused_moe/block_sparse_runcollisions as a separate, lower-priority cleanup — they are template reuse, not duplicate facades.fi_apistrings, then point decode benches/tests (CuTe + Cake) at the canonical facade.Rule for new kernels (immediate guidance)
New KDA work should do one of:
@flashinfer_apiwith a trace template + top-level/docs export);run_*stays inkda_kernelsas backend. State explicitly if a trace template is deliberately omitted, rather than defaulting to a bare decorator as the training ops do.RecurrentKDAPrefillWrapperdoes), or a paired forward/backward module (askda_training). Keep the implementation in the public module unless there is a stated reason for a private_*_implsplit.__all__/ top-level; state that clearly in the PR — the precedent set by feat(kda): add packed-input CuTe decode kernel #4417. With #4880 landing in0.7this becomes a formal track:@flashinfer_experimental_apifor an unstable public API, andflashinfer/experimental/plus@experimental_backendfor an opt-in backend reached through an explicitbackend=literal. Prefer that over widening a stable enum.Additionally:
backendliteral to one facade only (cf. feat(vibecuda_kda): add SM100/SM103 recurrent KDA prefill backend #4779). A new backend value is part of the shared contract.*_kda_decodenoun without stating why it cannot be abackend=value or a kwarg on an existing facade.kda_kernels.run_foo/kda_kernels.<submod>.run_foois the product API."Why this is backwards compatible
Compatibility argument, item by item
The compatibility baseline is
v0.6.18, not currentmain.flashinfer.kda_decode.recurrent_kda(...)calls keep working: the shim targets a decode dispatcher directly, so omittedbackendstays"cute-dsl"and the prefill classifier never intercepts a packed multi-token call. Its released two-valuebackendenum is only ever widened, never narrowed, and its default is never changed.fused_kda_decode/packed_kda_decode/RecurrentKDAPrefillWorkspace, and the released top-levelseq_order/prefill_workspace/ checkpoint kwargs, are unchanged.kda_kernelsnoun aliases andFLASHINFER_KDA_T1_FAST_PATHshipped, so both get a deprecation window rather than removal in place.RecurrentKDAPrefillWrapperare deliberately not held compatible withmain— they are unreleased, so step 3 revises them freely. "Backwards compatible" makes no claim aboutmain-only surface.flashinfer.recurrent_kda(...)onmainis already not behavior-identical tov0.6.18. It gained abackendparameter defaulting to"auto", which selects frozen Cake for eligibleT=1decode (feat(cake_kda): add native unbounded-softplus Kimi-Linear kernels #4535) and can raise where"cute-dsl"succeeds ([Bug]recurrent_kdabackend="auto"decode commits to Cake and raises instead of falling back to CuTe DSL #4935), and it now prefers CuTe DSL over Cake for prefill (feat(kda): add CuTe DSL recurrent prefill backend #4605/feat(kda): add SM120a CuTe DSL prefill backend #4633/fix(kda): fall back to Cake when CuTe DSL predates cutlass.experimental #4667). Because none of that shipped, steps 1–2 are free to reverse or ratify any of it; the only released constraint is that a defaulted call must keep working.Related
"auto"decode commits to Cake and raises instead of falling backbackend="cute-dsl"dense decode gathers and scatters the state pool in Python, so a negativessm_state_indicesentry wraps around and overwrites the last slot. Released inv0.6.18, silent, and independent of the unification plan — listed here only so the KDA decode defects stay in one place"cudnn"backend literal tokda.pyonly; step 0b@flashinfer_experimental_api,flashinfer.experimental; landing in0.7; the mechanism for the holding action above[RFC] Unified GDN Decode/Prefill API; the same problem class on the sibling opfeat(gdn): add unified decode API and deprecation shims; in-flight implementation of this exact pattern for GDNv0.6.17fortests/trace/test_template_registry.py::test_registration_module_inventory_is_complete, assertingExtra items in the right set: 'flashinfer.kda'. Already fixed and should be closed:flashinfer.kdais now in the static inventory (tests/trace/template_registry.py:74) and the test passes onmain(verified). It is evidence that introducing the second facade had knock-on cost, but it was never evidence for problem 2 — the assertion compares a hand-maintained module list against an AST scan, not label uniquenessImportErrorunder fix(kda): fall back to Cake when CuTe DSL predates cutlass.experimental #4667's version gate in environments with an older CuTe DSL; the first is an unguardedModuleNotFoundError: No module named 'cutlass.experimental'. Not caused by"auto"routing (that fallback is the mitigation), but the same theme:"auto"and the strict backends have materially different totalityRelationship to #4483
What #4483 got right, what it got wrong, and what is new since
#4483 raised this problem on 2026-08-12 and was closed on 2026-09-03 in favour of this issue. What changed in the three weeks between:
Resolved. Discoverability — #4483's fourth problem — is fixed:
backend=landed on the top-level facade in #4605 and is forwarded to decode, soflashinfer.recurrent_kda(..., backend="cake")now reaches the strict Cake path. Two pieces of #4483's guidance were also followed: #4417 merged internal-only, and #4378 was closed in favour of #4445's publicpacked_kda_decodenoun facade.Superseded by events. #4483 described a single "opposite backend policy" between the two facades. That framing no longer fits. The facades now share one
run_recurrent_kdaand differ in their defaults, and the top-level default is not a safe superset —"auto"decode commits to Cake and raises rather than falling back, which is #4935. Its prefill description has also inverted, but by events rather than by error: #4483 described"auto"prefill as preferring frozen FlashKDA, which was correct on 2026-08-12 and is still what releasedv0.6.18does (kda.py:260hardcodesbackend="cake"). CuTe-DSL prefill only arrived with #4605 on 08-20. Today"auto"prefers CuTe DSL with Cake as the fallback, and that fallback isnvidia-cutlass-dslversion-dependent (#4667).New since #4483. The
backend=parameter itself; the"auto"literal and the defect behind it; a training/backward family and a plan-and-run wrapper class (six top-level symbols, all decorated-but-untraced or classes, plus a private_kda_training_implsplit); an SM120a prefill tranche that added threekda_kernelsexports and a new subpackage; two more docs pages; and concrete fragmentation across seven open PRs, with #4709 as a counterexample that keeps both facades in step. Critically, none of that has shipped — see Release status — which is what makes the plan above cheap now and expensive after0.7.Unchanged and still reproducing. The duplicate
recurrent_kdaname and shared trace label; the top-level facade bypassing the decode wrapper; thestage:decodetag on a template that dispatches prefill; missingbackend/seq_order/prefill_workspacetrace fields;kda_decode.rstautosummarisingrecurrent_kda; thekda_kernelsnoun aliases; decode-only fi_trace coverage; and the prefill/decode test import split.