v1.28.0 release

cuDNN Frontend v1.28.0 Release Notes

cuDNN Frontend v1.28.0 is the recommended version for cuDNN 9.25.1 and later releases.

New: cudnn.fla — a drop-in accelerator for flash-linear-attention 🚀 🚀

cudnn.fla (#596) monkeypatches the flash-linear-attention ops that cuDNN can serve onto cuDNN's Blackwell (SM100) kernels, with a transparent fallback to FLA everywhere else, so results never change:

import cudnn.fla
cudnn.fla.accelerate_fla()   # before importing FLA layers/models
import fla                   # GatedDeltaNet / KDA now run on cuDNN where supported

• Gated DeltaNet (chunk_gated_delta_rule) — the GDN convention (log-space decay g, post-sigmoid beta, GVA where HV > H) mapped onto cuDNN's native op, reproducing the fused-layer knobs use_gate_in_kernel, use_beta_sigmoid_in_kernel, and use_qk_l2norm_in_kernel.
• KDA (Kimi Delta Attention, chunk_kda) — channel-wise gate plus scalar beta, l2norm forward and backward through cuDNN. BF16 only; FP16 declines and falls back.
• GatedMLP (#686) — an opt-in adapter (accelerate_fla(targets="gated_mlp")) backed by cudnn.gemm.ops.swiglu_mlp. The patch registry is target-selective, incremental, idempotent, and independently restorable via restore_fla(targets=...) / is_accelerated(target).

Configurations cuDNN cannot serve raise cudnnGraphNotSupportedError / NotImplementedError and fall back to FLA — never a wrong answer. Correctness is pinned by test/python/linear_attention/test_fla_compat.py, which requires cuDNN to match FLA within FLA's own BF16 noise on the output and every gradient.

Underneath, the linear-attention stack gained KDA and GDN-2 backward support (#556), a safe beta guard for GDN-2 (#722), packed-QKV views for native GDN (#685), state-layout and convention alignment with FLA/FlashInfer plus a context/IMA fix (#644), and successive CPU-overhead and instruction-cache/numerics passes on the FROST linear-attention kernels (#616, #708, #759). See docs/fe-oss-apis/fla.md.

New: JAX support across the CuTeDSL GEMM APIs 🚀 🚀

The GEMM CuTeDSL APIs are now type-erased (#529): every API under python/cudnn/gemm/cutedsl/ accepts JAX arrays alongside torch tensors, and the modules import and resolve their public symbols without torch installed — torch is imported only when torch tensors or dtypes are passed, and JAX only when JAX arrays are.

On top of that, cudnn.jax.call (#553) wraps CuTeDSL's native JAX integration (cutlass.jax.cutlass_call) and gives every JAX-reachable GEMM API a jax.jit entry point — gemm_amax, gemm_swiglu (including blockscaled MXFP8), gemm_srelu, gemm_dsrelu, gemm_proj_rope_mxfp8 (both BF16 and MXFP8 input paths), and the grouped and discrete-grouped families in their pointer-array modes. APIs without a JAX data path raise a clear error rather than failing obscurely. JAX outputs the kernel already writes are no longer zero-initialized first (#631).

New: First-class cudnn.Handle 🚀 🚀

cudnn.create_handle() now returns a Handle object that owns {backend handle, device, stream} instead of a bare int (#612). The per-handle state that had accreted as module-global side tables and per-engine device queries — the stream cache, and the three parallel device stacks used by the backend handle, pygraph, and FROST — unify behind Handle.stream and Handle.device. This matters because the Python engines (FROST, CuTeDSL, linear attention) need a device and a stream, not a cudnnHandle_t.

Backward compatibility is transparent for normal use: every handle-taking API (execute, set_stream, get_stream, destroy_handle, all graph methods) is Handle-aware, extracting the backend handle explicitly at each named handoff. Design notes and a full call-site inventory are in docs/handle_first_class_design.md.

New: GNN simple aggregation 🚀 🚀

from cudnn.gnn import CscGraph, agg_simple

cudnn.gnn.agg_simple (#647) exposes the cuDNN GNN AggSimple backend as a PyTorch custom operator with autograd, fake-tensor, and torch.compile support, handling graph validation and backend invocation so callers never touch the low-level GNN structures. Requires cuDNN 9.26 or newer and compute capability 8.0+; not supported on Windows. See docs/operations/gnn/agg_simple.md.

New: FROST SDPA on Ampere, Ada/Blackwell-consumer, and Rubin 🚀 🚀

The FROST engine family introduced in v1.27.0 now spans every architecture the frontend targets.

• SM80 (A100) — CuTe-DSL prefill (generic and d=256) and backward (generic and a d=64 fast path, roughly 2× on plain dense MHA) engines plus standalone cudnn.sdpa adapters (#493), later ported to plan-time compilation with TemplateParams kernels and sym_int THD extents (#689). Previously the manifest had no engine below SM100.
• SM120 — a forward per-tensor FP8 (E4M3) engine (#509) grown through several rounds of features and tuning (#531, #562, #595), MLA d_qk=192 / d_v=128 support (#507), and a backward engine (#486) extended to all head sizes ≤256 including d192/d256, non-compact layouts and deterministic dQ (#533), sliding-window attention (#505), GQA/MQA, padding masks, right-band widening and sink-token gradients (#557), deterministic 2-kernel mode and dBias (#707), and native service of declared strided layouts (#666).
• SM107 (Rubin) — per-tensor FP8 prefill via a sibling kernel (#576), with the has_lse specialization and a static SMEM guard (#579), a fused LDTM row-max and row-sum-in-MMA epilogue (#580), and the softmax_precision knob axis lit up with an F16x2 exponent on the d128 sibling (#651).
• MLA backward — sdpa_bwd gains MLA support (#643).

Head-dim envelopes and engine identity. Per-tensor FP8 now serves the dense head-dim ENVELOPE through the same TMA zero-padding path the F16/BF16 flavors use, and the engine table collapses to one engine per architecture × dtype family — head dims became a lowering concern (kernel-flavor selection) rather than an engine identity (#587).

Ragged / THD, with zero host reads. THD execute on SM100 (#606) and SM120 (#608) now performs no device-to-host reads at all — no .tolist() syncs, no host cumsum, no pageable H2D — building its metadata on device against a plan-time envelope grid, which makes the path CUDA-graph capturable (issue #552, with #543 binding host prep to the launch stream and plan-time-only THD compile keys). The FP8/MXFP8 SM100/SM107 engines were moved onto the same envelope design (#648) and the legacy pre-envelope THD leg removed (#622). Supporting work: native THD declared-stride support in the SM100/SM120 F16 forward kernels (#526), the cu_seq_len prefix-sum length form (#522), ragged stats on SM100 (#512) and SM120 (#508), and ragged S_kv tails served on the F16 rows via synthesized padding (#581).

Masking, splitting, and heuristics. Forward heuristics can now recommend the same engine several times under different complete knob assignments, which makes split-KV graph-reachable for the first time (#692); recommend() is a pure, backend-blind entry point that autotuners can call with hand-built graph facts. Split-KV also landed for the SM100/SM120 prefill kernels (#658), the KV split the heuristic chose now runs on the true cluster shape (#720), and pack_gqa is supported (#709). On masking, SM100 gained causal right-band widening with per-sequence THD bottom-right diagonals (#485), bottom-right diagonal plus sliding window (#584) — after which the bottom_right_with_swa notch was retired because every row serves BR + SWA (#623) — and full causal mask support (#498).

Other FROST SDPA work. SM100 MXFP8 for d192/d128 (#661); dense LSE written directly to non-contiguous, dense-compatible layouts (#712); an execute path made async where it can be, no longer re-deriving build-time facts (#570); FP8 scales folded in-kernel with a baked 2⁴ P-cast bias, removing Scale_S from below the graph (#619); the Amax_S output dropped from the FP8 kernels (#602); a has_lse specialization for the FP8/MXFP8 SM100 flavors (#574); and a strict LSE/sink/seq-lens execute contract with no torch.empty in execute (#484).

Removal: the legacy standalone SM100 d=256 forward and backward stacks, along with the cudnn::sdpa_{fwd,bwd}_d256 experimental torch ops; SM80 forward moved onto the same SdpaFwdDsl adapter path SM100/SM120 use, so one lowering function drives every forward cell (#682). d=256 remains available through the graph API on backend engines.

Updates to Graph API 🚀 🚀

SDPA

• max_total_seq_len_q / max_total_seq_len_kv on the forward node (#740). sdpa_backward has accepted these since cuDNN 9.6; the forward node never did, so a ragged graph could not express its packed token total and the FROST forward path had to infer a loose upper bound from the bound buffers' element span. A loose bound is memory-safe but not benign — masked rows are still multiplied, so an over-allocated, unwritten tail poisons whole tiles through 0 * NaN. Every framework already holds this number (q.shape[0] in vLLM, SGLang, TransformerEngine, Megatron-Core, PyTorch, FlashInfer); it can now be declared.
• Stats must be FP32 (#690). The SDPA kernels always compute and store logsumexp as FP32. A graph that declared Stats with a narrower dtype — explicitly, or implicitly by leaving it unset with a non-FP32 io_data_type — built and executed fine, and the kernel then wrote FP32 rows past the end of the caller's buffer, surfacing as silent corruption of adjacent allocations, illegal memory accesses, or driver launch failures. Stats is now set to FLOAT at creation, an unset dtype defaults to FLOAT, and a narrower declared dtype is rejected.
• Non-BHSD softmax-stats strides are rejected before cuDNN 9.26 (#304, NVBug 6057616). The SM80 and SM100-dBias backward kernels in cuDNN < 9.26 ignore the declared LSE strides and address the tensor as packed BHSD, silently producing wrong gradients for any other layout. The forward graph now fails at construction rather than the backward failing silently at runtime; the pre-9.26 packed-BHSD check was moved to post_validate_node (#642).
• Capabilities.bottom_right_padded_seq_q was retiled (#683).

Serialization and plan management

• Handle-less plan deserialization (#544) — a Graph::deserialize(blob) overload (and pygraph.deserialize(blob)) rehydrates a serialized execution plan from a DeviceProperties descriptor instead of a cudnnHandle_t, enabling ahead-of-time compilation: build and serialize a plan on a GPU node, deserialize it later where no CUDA context or cuDNN handle exists. Requires cuDNN ≥ 9.8 at compile and runtime; the API compiles on older headers and returns a runtime error.
• Tensor_attributes::alignment is now serialized (#564).
• CUDNN_KNOB_TYPE_TILE_CGA is mapped (#729). An engine reporting a knob the mapping did not carry returned it as NOT_SET, and feeding that back through create_execution_plan() failed for every knob combination on that engine — making the engine impossible to drive through the explicit-plan API at all. KnobType_t::TILE_CGA is added and mapped in both directions, and exposed to the Python bindings.

Python dispatch

• One place ranks plans (#528). propose_plans, the Router, and heuristics_sort are deleted. An engine cannot rank plans — it sees neither its siblings nor the backend's entries — and all four in-tree propose_plans were the base class's default copied verbatim. create_execution_plans() now goes straight to heuristics.rank(...), which delegates to each family's recommend().
• The variant pack is normalized once (#547), into a C type that is also the DLPack producer, so the cuDNN backend and every Python engine see the same object below graph.execute(). This closes three cases where the same public call answered differently depending on which plan the heuristics happened to pick — bare device addresses, override_shapes on a FROST plan, and related identity-dependent behavior.
• Python engines are grouped by families, with assorted cleanups and fixes (#502).

Operations

• FFT causal conv1d bindings are back (#542). Reintroduces the support merged in #437 and reverted in #479, now that the cuDNN 9.26 release branch carries the required backend declarations: dynamically loaded pybind shims for the medium and long paths (including the long-path workspace/reserve-space size query), cudnn.ops.fft_causal_conv1d(x, weight) following cuhyena's medium/long selection, padding, trimming and autograd behavior, preserved long-forward reserve space for the matching backward, C++ samples, notebooks, and docs. SM107 causal conv1d tests are skipped before cuDNN 9.26 (#632).

Open-Source Kernels 🚀 🚀

GEMM and MoE fusions

• cudnn.gemm.ops.swiglu_mlp (#609) — a dense BF16 autograd op for out = (silu(x @ Wg.T) * (x @ Wu.T)) @ Wd.T on SM100. The forward gate/up GEMMs, SiLU, and multiply run as one FORT-native runtime-fusion kernel that also emits gate and up, avoiding two recompute GEMMs in training; the backward fuses dh = dout @ Wd with the two-output dSwiGLU epilogue in one FROST kernel, keeping dh on chip. Unsupported layouts, architectures, and missing optional dependencies decline to nvjet plus pointwise.
• Rubin (SM107) MoE fusions — a fused block-scaled grouped GEMM + SwiGLU + RHT + NVFP4 quantization kernel producing both row-major and column-major quantized outputs with swizzled block scale factors (#637), built on new E5M3 scale-factor support for fused GEMM on Rubin (#545); the GLU+RHT+quant kernel now emits column-wise RHT in ragged tensor layout (#693); FROST GEMM gained Rubin kernel pipelines (#593); and single-group runtime offsets were added to Rubin grouped dGLU (#590).
• Grouped SiTU-GLU activation (#645), with the SiTU-GLU Hadamard API contracts fixed (#670).
• Deterministic dprob for grouped GEMM dsrelu (#521) and FP32 row-scaled FP4 grouped GEMM (#461).
• Host overhead cut roughly 3–4× on the CuTeDSL grouped-GEMM wrappers (#627) by memoizing work that is a pure function of values the memo keys already hold — no check skipped. On B200 at 2048×2048×2048 with 8 experts: grouped_gemm_wrapper_sm100 118.1 → 39.6 µs, grouped_gemm_glu_wrapper_sm100 161.1 → 40.6 µs, grouped_gemm_dglu_wrapper_sm100 185.0 → 51.1 µs (kernel time 18.4 µs).
• The public GEMM path now works for the flavors it claims (#558), with non-MoE, workspace-free flavors lowered off one build-time recipe (#559).
• Kernel-level throughput work: mixed CGA for the FROST GEMM kernels (#617), multiple tcgen05 instructions in the M dimension (#566), multi-UTCMMA optimization (#511), the LPT_L2 scheduler (#585), TMA-STG with multi-GEMM (#721), M-major output refactorization (#735), and Rubin MoE GEMM optimization passes (#667, #697, #700).
• SM120 matmul support (#719).
• The NVRTC dependency for occupancy queries was removed (#494), and the elect_one compilation hint optimized (#504).

DSA (DeepSeek Sparse Attention)

• Opt-in SM100 sparse indexer backward v2 (#640) — backend="sm100_v2", a drop-in for the SM100 sparse indexer backward that is 1.16–1.92× faster on the GEMM stage (1.92× at topk=128, 1.31× at topk=1024, 1.17× at topk=2048; 1.31× end-to-end through the public wrapper at S=8192/topk=1024). A two-term BF16 hi/lo expansion of A = g·w additionally keeps d_index_k FP32-accurate at ~no cost for consumers that keep index_k in FP32. Scope is SM100 exactly, H=64, D=128, topk ∈ [128, 2048] in multiples of 128, sm_scale > 0, request-or-fail with no silent fallback; the default backend is untouched.
• The SM100 sparse-attention backward kernel was sped up (#684), and an SM100 H16 backward specialization added for H=16, d_qk=576, d_v=512 BF16 (#664).
• indexer_backward now validates the output and plan signature before kernel 1 on the default SM100/SM90 backends (#572), and range_constexpr was restored in eight kernel_gemm epilogue loops (#549).

Block-sparse attention (BSA)

• Sage FP8 forward for Blackwell (#475) — a new public block_sparse_attention_fp8_forward API that quantizes contiguous BF16 BHSD inputs to FP8 E4M3 internally using the Sage recipe and returns contiguous BF16, gated on CUTLASS DSL 4.6.1 at runtime. Covers SM100/SM103 (blk64, with automatic split-KV selection) and a dedicated SM120 kernel with sequence tails, fixed or variable sparse counts, and batched block_sizes layouts. Persistent CLC scheduling now works together with split-KV: the scheduler's work-tile mapping explicitly encodes and decodes the split dimension.

CSA (Compressor)

• Enabled on SM100+ GPUs generally (#641), with the validated ratio envelope extended to ratio=2 (#710) — ratio ∈ {2, 4, 128} with coff ∈ {1, 2}, the configuration used in production training for the model family this operation serves. No kernel changes; the previous gate encoded validation scope, not a kernel limitation. Review-response fixups for the ratio=128 kernels landed in #452.

Toolchain

• The FROST engines want nvidia-cutlass-dsl ≥ 4.7.0 and check the version at support time, declining rather than failing when the installed DSL is older; the package itself is deliberately not pinned to that floor so it stays compatible with consumers holding the DSL back. The packed-FP4 wgrad layout workaround is now gated on cutlass-dsl < 4.8 (#764).

Tooling, CI, and Build ✨✨

• Style checks moved to GitHub Actions (#741) — pre-commit now runs as a GitHub Actions job.
• Merge requirements are enforced (#736, #739) — a merge-requirements job fails while a PR has no Milestone or is not on any Project board, with a matching PR-template checkbox. Bot-authored PRs and PRs labeled cat-routine-update are exempt, and a failed Projects lookup reports a clear error rather than a false pass.
• Stale-PR automation (#714).
• setup.py defaults to a parallel extension build (#565), and a -Werror unused-parameter build break in init_gnn_submodule was fixed (#728).
• The large-tensor fuzzer's plan fallback and failure output were improved (#583), on top of a new large-tensor convolution fuzzer (#401).

Samples, Benchmarks, and Tests 📊

• Attention inference benchmark suite (#636) — benchmark/attention_inference/ measures attention as served, in two phases: context (TFLOPS; full prefill and chunked prefill of 512/1024-token chunks against 64k/128k caches, bottom-right causal) and generation (GB/s and % of memory SOL; q_tokens = 1 + MTP for MTP 0–3 against a 128k cache). Two backends are swept and charted — cudnn on native backend engines, and cudnn_oss planned with heur_mode.OPENSOURCE plus CUDNN_FRONTEND_ENABLE_FROST_ENGINES=1 so only the frontend's open-source engines may serve it, recording the winning plan per case.
• Reproducible end-to-end benchmarks (#687) — a Qwen3.8 2³ layer-period benchmark with an eight-treatment Williams design, paired effects, and Shapley attribution over independent GDN, SwiGLU MLP, and d256 full-attention treatments, plus a pinned Qwen-Image BF16 transformer proxy built on Diffusers' real QwenImageTransformer2DModel. These exercise the merged production paths: packed-QKV GDN (#685), the cudnn.fla GatedMLP shim (#686), and the backend-only public d256 SDPA after #682.
• The attention benchmarks were restructured (#657) and the benchmark scripts unified (#655); the SDPA benchmark gained the cudnn_oss (FROST) backend with unified sustained-clock SOL and peak lines (#597), FA4 auto split-KV enabled by default (#607), a Qwen3-VL vision-encoder (ViT) config with GB300 results (#598, #629), SM-clock sampling on the GPU the benchmark actually runs on (#699), and an Ampere (sm80) row in the peak-MMA table (#715). GDN benchmarking was added in #501.
• Benchmarking artifacts refreshed against backend 9.26.0.17 and 9.26.0.33 (#657, #665, #681, #694, #711).
• Tests: ragged per-tensor token-stride gaps are fuzzed in test_mhas_v2 (#516) and sink tokens are drawn in the ragged backward suites (#630); DSA comparisons use only the effective top-k slice (#672); render-only tests were removed (#688); samples skip an unsupported case and drop an invalid cudaGraphDestroy (#734); the 02_low_level_api notebook uses an FP32 Stats tensor (#701).
• CausalConv1d sample links and the backends architecture fence were fixed (#518).

Bug Fixes 🐛

SDPA

• Ragged backward workspace under-allocation with head-major stats (#462). A ragged sdpa_backward graph that set max_total_seq_len_q silently returned all-zero dQ/dK/dV when Stats was head-major — the forward O and stats were correct and no error was raised, so this corrupted training without ever surfacing. Head-major [h, total_q] is FlashAttention's and PyTorch varlen's softmax_lse layout, i.e. exactly the integrations that would reach for max_total_seq_len in the first place. Present unchanged in 9.22 through 9.26 on both sm90 and sm100.
• Fixed an SM100 FROST zero-KV cluster deadlock and enabled zero-length seq_len_kv sweeps (#575).
• Fixed a latent THD O-descriptor row stride on the SM100 FP8 path (#577).
• Fixed SM100 D192 E5M2 attention-sink accuracy (#691).

Python and device handling

• A CUDA context is now bound on the calling thread — the right one, and on both sides of the boundary (#626). A driver-API launch reads the calling thread's context stack, and a thread that has done no CUDA work has nothing on it; ensure_current_context also returned as soon as any context was current, so a thread bound to another GPU's context kept it. It now resolves the target context instead of accepting the incumbent. The missing ensure_current_context imports introduced with #612 were added in #638.
• The oversized-SMEM ceiling is requested by ordinal when the binding cannot name the device (#635).
• DLPack is told the frontend is only reading, so probe() survives a graph capture (#506).

Kernels

• Three Rubin (sm_107) compile and launch failures fixed (#696), including all 26 FP8-path test_grouped_gemm_quant tests failing with CUDA_LAUNCH_INVALID_CONFIG: _compute_stages picked its A/B pipeline stage count from a byte-sum model that ignored the 1024-byte alignment padding of SharedStorage, so the launch requested 335,872 bytes of dynamic SMEM against an sm_107a cap of 334,848.
• The TMEM accumulator is read before the overlapping stage is released, in grouped GEMM (#654) and in SM107 grouped GEMM quant (#673).
• Fixed the dGLU xfilter math (#568) and a missing attribute in Rubin grouped dGLU (#625).
• Fixed missing engine class imports in the linear-attention tests (#513) and a CI breakage from a missing is_exclusive guard (#605).

Acknowledgements 🙏

Thanks to everyone who contributed to this release:

@Adnios, @adshen, @Anerudhan, @Aneureka, @brandonfzhang, @Butterfingrz, @dextero, @egilliam-nv, @filaretov, @GarlGuo, @harryzhou2000, @hwanseoc, @hxbai, @icavan, @jhjpark, @jiayus-nvidia, @kainzhong, @msalasooNV, @pmdavies-nv, @sraman-rgb, @timmoon10, @tingyu66, @tp5uiuc, @vasunvidia, @vedaanta, @YangXu1990uiuc, @yanqinz2, @yanzhuo607, @yeliu-oss, @yihuawei, @ZhiyuLi-Nvidia, @zianglih, @zkyue, Ali Hassani, Shreya Gaur, huangyitong.hyt@alibaba-inc.com and wenting.swt@alibaba-inc.com.