cuDNN Frontend v1.29.0 Release Notes

cuDNN Frontend v1.29.0 is the recommended version for cuDNN 9.26 and later releases.

New: HSTU attention 🚀 🚀

import cudnn
from cudnn.hstu_attention import hstu_attn_varlen_func

HSTU (Hierarchical Sequential Transduction Unit) attention arrives as a complete CuTe DSL kernel family for Blackwell (#487) — packed variable-length forward and backward, FP16/BF16, head dimensions 64, 128, and 256, with full, causal, local, and arbitrary masks, paged-KV forward, and strided or preallocated gradient outputs. HSTU replaces softmax with a SiLU score transformation and derives block-sparse metadata automatically. Explicit-stream execution is allocation- and lifetime-safe, cross-device PyTorch streams are rejected, and output overlap is validated against all read-only metadata.

Built out over the release:

• D32 head dimension (#885) and LMSD — a BF16 LayerNorm-Multiply-SiLU-Dropout forward and backward pair for hidden dimensions below 1024 divisible by 8 (#946), with independently strided X/U inputs, optional SiLU, dropout, concatenated U/X outputs, and optional dWeight. Launch and vector configuration are chosen from the hidden dimension and the runtime device SM count, and changing a row stride does not trigger recompilation. Independent backward output strides followed in #962.
• qlen=1 causal and local attention optimized on SM100, SM103, and SM107 (#900), after a layout refactor that cleaned up the qlen=1 kernels (#927).
• Fixes: TMA store synchronization (#784), block-sparse test input ordering before the cross-stream build (#938), and the removal of the cross-tensor storage overlap checks (#1012).

New: DSA sparse attention forward — the loop closes 🚀 🚀

cudnn.DSA.SparseAttentionForward and sparse_attention_forward_wrapper (#569) add the SM100 sparse forward path, so the DeepSeek Sparse Attention forward/backward workflow now completes inside the frontend-only CuTe DSL API instead of requiring an external FlashMLA forward. H64 D512/D576 and the H128 D512 small-top-k prefill specialization are covered, including arbitrary logical top-k lengths, invalid/out-of-bounds and duplicate indices, per-query lengths, attention sinks, and an optional indexer LSE. Supported on SM100-family capabilities 10.0, 10.3, and 10.7; decode, split-KV, regular H128, SM90, and FP8 cache paths are not included.

On the backward side:

• Deterministic SM100 backward with a bounded-wave dKV reduction (#810).
• Two-CTA specialization for BF16 H128/D512 — 1.10–1.16× faster (#817), an H32/D576 specialization (#760), and optimized indexer-backward gradient kernels (#730).
• Gather4 and score kernels optimized, with sparse attention documented (#957).
• Fixes: three SM90 top-k and attention-sink boundary failures (#785), invalid sparse rows on the SM100 backward (#877), backward boundaries alongside relaxed SM100 BSA FP8 shape limits (#926), and deterministic handling of infinite sinks (#940).

New: Flex Attention 🚀 🚀

import cudnn
mask_plan = cudnn.create_mask_plan(...)
out = cudnn.flex_attn_func(q, k, v, mask_plan)

The experimental cudnn.flex_attention namespace (#775) brings CuTe DSL forward and backward kernels for SM90, SM100, and SM103, covering fixed-length and variable-length MHA/GQA workloads, with compact arbitrary-mask planning, scheduling, runtime compilation and cache support, autograd integration, and lazy top-level exports. The port also syncs the CUTLASS DSL 4.6.0/4.6.1 bulk-copy election fix that otherwise deadlocks the SM90 backward. An SM100 2-CTA mask-slot synchronization bug was fixed in #993.

New: torch.sdpa runs on the cuDNN Python API 🚀 🚀

torch.sdpa is now served end to end — forward and backward, dense and varlen — by the cuDNN Python API (#554).

The "CUDNN" provider (python/cudnn/torch/) registers with torch.nn.attention's flash-impl registry (PyTorch 2.13+, the same mechanism FA3/FA4 use) and overrides the CUDA kernels of aten::_scaled_dot_product_cudnn_attention{,_backward}, so F.scaled_dot_product_attention under sdpa_kernel([CUDNN_ATTENTION]) and torch.nn.attention.varlen.varlen_attn run on pygraph plus the engine Router. Registration is passive; activation stays explicit.

Alongside it, cudnn::sdpa_bwd gained dense BHSD backward — it previously served only packed THD and raised NotImplementedError on dense. With that in place the provider's dense backward stops falling back, closing the last C++ hop in a dense training step. Three defects were fixed in the same change: 2·B blocking device-to-host syncs per varlen backward (the host loop is replaced by the device-side thd_lse_to_padded() conversion, which makes the path capturable), dense operands never being normalized to the innermost-dense / 16B-aligned-base contract, and the dense autograd path dropping is_deterministic so use_deterministic_algorithms(True) still built a non-deterministic backward.

The experimental torch op was replaced by cudnn::sdpa_fwd / cudnn::sdpa_bwd (#780), which then gained window_right and asymmetric band service (#901).

New: Paged KV caches on the FROST SDPA engine 🚀 🚀

Paged KV caches are served by the existing SDPA graph API and the FROST SM100 engine — no new entry point (#964). A graph built with cuDNN's own paged-cache contract (paged_attention_k_table / paged_attention_v_table plus paged_attention_max_seq_len_kv) now lowers onto the d128 f16/bf16 kernel's new PAGED_KV specialization. FlashInfer's cudnn_batch_decode_with_kv_cache already builds exactly this graph, so with CUDNN_FRONTEND_ENABLE_FROST_ENGINES=1 vLLM and SGLang decode reach the FROST kernel with zero integration work.

• HND ([num_pages, H_kv, page_size, D]) and NHD ([num_pages, page_size, H_kv, D]) page layouts — the layout is only the pool's strides, read off the bound strides, with no new TemplateParams field.
• d128 and d256 f16/bf16 flavors; d=64 rides the d128 envelope and d=192/192 the d256 one.
• THD (ragged) queries over the paged cache — chunked prefill on the same engine.
• Per-batch KV lengths read in-kernel; num_pages / max_pages never enter the compile key. graph.execute runs clean under torch.cuda.set_sync_debug_mode("error"), and the adapter path captures under torch.cuda.graph with seq_len_kv contents changing between replays.
• KV split + combine is proposed for paged graphs — the padded exclusion is lifted, since paged graphs are padded by construction.

FP8/MXFP8 pools, sinks, unaligned page sizes, d > 256, a missing padding mask, and packed block tables are declined at plan time.

FROST SDPA: the large-head-dim band, Rubin, and Ampere 🚀 🚀

SM100 d=512 backward (#887). sdpa_bwd_sm100 serves the band d ∈ (256, 512] on Blackwell; before this, every sdpa_backward() there fell through to the cuDNN backend, which does not serve it. It is a three-stage chain, not one fused kernel — a fused d=512 backward needs 512 TMEM columns for dV and 512 more for dK against 512 per CTA, so S and dS go to a GMEM workspace and the gradients become three batched GEMMs (do_dot → bprop_d512_f16_sm100 → bprop_matmul_sm100). Stage 2 forks the forward d512 kernel's cga4x1 role split. Two user-visible consequences are documented in the support matrix: the workspace is 2·B·H_chunk·S_q·S_kv·2 B (the host loops over head chunks to hold it under 4 GiB), and the band is envelope-served, so d=264 pays d=512's MMA cost. It covers GQA/MQA, causal (top-left and bottom-right), SWA, right-band widening, arbitrary S_q/S_kv, and BSHD plus arbitrary dense stride order. THD/ragged support followed in #898, GQA/MQA under THD in #922 (retiring both THD conjunction flags), and a sXfer bank-conflict fix in #941.

Quantized and large head dims. Per-tensor FP8 d512 forward on SM100 (#845), SM100 D512/D512 MXFP8 prefill forward (#989), a specialized SM120 D512 F16 prefill kernel (#991), SM100 D256 FP8 and MXFP8 kernels (#860), the SM100 d=256 MXFP8 backward engine sdpa_bwd_sm100_mxfp8 (#904), a specialized SM120 f16 d256 kernel (#930), and completed SM100 D192/D128 feature support with performance tuning (#841). Split-KV partials are always stored in fp32 on SM100, and KV split is allowed with a quantized output (#891).

Rubin (SM107) prefill forward (#954) — nine new kernels joining the already-shipped d128 FP8 sibling behind three opt-in engine rows (sdpa_fwd_prefill_sm107, sdpa_fwd_prefill_sm107_fp8 widened to d256/d512, and sdpa_fwd_prefill_sm107_mxfp8), spanning f16/bf16, per-tensor FP8, and MXFP8 across d128, d192×d128, d256, and d512. Rubin gets its own config_sm107.py because it is a different kernel lineage, not the Blackwell kernels recompiled: routing the SM107 d256 FP8 kernel through the Blackwell make_cfg_d256 produced a scheduler mbarrier init count of 15 against 11 actual arrivers — an unreachable barrier that hung at every shape, down to B=1 H=1 S=128, with no fault and nothing for a sanitizer to find. Follow-ups added DSv3 quantized flavors, THD on every f16/FP8 flavor and arch-grouped forward kernels (#974), SCHED_LPT on the fp8 row (#850), SCHED_LPT honored with STAGES_KV restored as a knob (#1001), and SCHED_LPT claimed at d256 / d192×128 with heuristics ranking from the flavor's domain (#1020).

Ampere (SM80) backward. The q-loop is now bounded by the sliding window — 6–12× on gpt_oss-style SWA — with a window-aware deterministic relay (#866). THD plus attention sinks and THD plus deterministic dQ landed in #867, native strided-LSE reads and a THD max_s_kv grid hint in #766, the port onto SdpaBwdDsl + TemplateParams with sym_int THD extents in #765, workspace carving that removes per-execute allocation on the engine paths in #716, and conflict-free dQ staging with an L2-grouped causal grid order in #948.

THD, scheduling, and plan selection. THD LPT remap (#717), THD persistent grid (#848), split-KV ports (#768), native strides (#795), the CLC scheduler aligned with the canonical CUTLASS pattern (#894), and python-native validate() per engine family (#869), which defers the C++ lowering while a Python engine is still a candidate and fixes the issues that exposed on the SM80 backward and SM100 FP8. FROST GEMM and SDPA kernel symbols are now prefixed with cudnn (#854).

Linear attention and causal convolution 🚀 🚀

• kda_cake (#912) — a third engine for kimi_delta_attention, hosting the CAKE-generated recurrent KDA training kernels. The kernel bodies are vendored byte-for-byte from FlashInfer (with UPSTREAM.md and SHA256SUMS), compiled with NVRTC at first use and launched through the driver API, with an on-disk cubin cache keyed by source digest, options, and NVRTC version. The engine is opt-in — check_support declines unless CUDNN_FRONTEND_ENABLE_FROST_ENGINES=1 is set. Training kernels ship from cuDNN Frontend and inference kernels from FlashInfer; hosting CAKE here next to kda_frost and kda_cutile keeps one NVIDIA training path for the operator and gives it autograd, the cudnn.fla drop-in, and the JAX path for free.
• Native causal conv1d full-sequence and decode operations (#799).
• GDP support for FROST LA (#742), context-parallel primitives plus improved low B·H performance (#969), and gate + beta in F16 with PDL enabled (#819).
• Fixes: a GDN regression alongside a FROST LA cleanup (#890), a split-K heuristic fix (#790), NaN from the cuTile engine at headdim=32 for GDN (#994), and stale TMA descriptor flags for short varlen tails (#1015, #1013).

Open-Source Kernels — GEMM, MoE, and quantization 🚀 🚀

• Mixed-input and mixed-precision GEMM: MXFP4 × MXFP8 (#970), FP8 × MXFP8/NVFP4 (#996), and 2×FP8 regular GEMM (#1022), plus swap-AB (#1019).
• Two-kernel split-K for SM100 (#856), extended with epilogue fusion and performance tuning (#942) and then to SM120, SM100 block-scale, and SM103 block-scale with a swept split-K configuration (#963).
• SM120 matmul: block-scaled matmul support (#945) and a configurable path covering all m/n/k-major combinations (#816).
• MoE: a dynamic scheduler (#749), segmented row-scale quantization (#812), grouped-matmul performance work (#788), a model-specific SwiGLU MoE API (#899), and a fix that drains the final scheduler broadcast (#813).
• Kimi K3 SiTU dense MLP support (#807).
• Epilogues: an optional tanh_clamp_scale soft clamp for the squared-ReLU grouped GEMM epilogues, forward and backward (#858), with its instruction count subsequently cut (#906); row/column-wise quantize optimization (#833); epilogue tile-size selection (#849); and 1-CTA FROST plans preferred for quantizing epilogues (#808).
• Heuristics: real-M scoring with a densified tile scan that never pads the cluster grid (#844), a 64-tall tile family for M-starved shapes with measured cluster tie-breaks (#781), and block-scale early pairing gated on an exact M rather than a MoE group average (#871).
• GEMM code refactor (#776).
• Triton NVFP4 attention QAT backward (#778).
• BSA: SM100 blk64 partial KV tails served without padding (#861).

Graph API, Python API, and packaging ✨✨

• CuTeDSL is now a required dependency, with the floor raised to 4.6.2, and requires-python raised to >= 3.10 (#921). The OSS kernel engines (FROST SDPA, GEMM fusions, sparse and linear attention) are Python kernels that JIT through CuTeDSL — a headline part of the package rather than an add-on — so the runtime they need installs by default. The DSL version is additionally gated at runtime so an older DSL declines rather than fails (#917), with follow-through on the SM80 SDPA gates, a prerelease-aware floor, and a lazy-import message (#919).
• Kernel cache: thread-safety guarantees (#815) and revision() / size() accessors (#909).
• SDPA validation: Bias dim and stride are validated like every other SDPA I/O tensor (#823); the forward now accepts any Stats output layout, dropping the pre-9.26 packed-BHSD guard (#1023).
• BackendDescriptor's constructor avoids copies (#761).
• The obsolete NumPy upper bound was removed (#1005).

Documentation 📖

• The cuDNN documentation now lives in this repository and builds with Fern (#1006) — the published docs build from the same tree as the code they document. The nine loose docs/*.md files moved into docs/utilities/, the RST-authored guides were converted and brought over, and fern check and publish workflows were added. Markdown across the tree was made MDX-safe (#968), and the docs site base path moved to /cudnn (#1016).
• A FROST SDPA support matrix was added, with a rule to keep it current (#865); O-stride semantics in the fp16_fwd sample were clarified (#756); and AGENTS.md codified the THD Stats packing rule, the editable-install gotcha, and a PR review checklist (#843).

Tooling, CI, Benchmarks, and Tests 📊

• Crashing and hanging tests are isolated so one bad test cannot poison the run (#639), with the per-test deadline raised to 1500 s (#911).
• A registry-driven SDPA fuzz framework — test/python/sdpa/suites, covering context, generation, and bprop across dtypes and model presets (#846). test_mhas_v2 gained flash-style references using roughly 16× less memory, sequences up to 8192, and more L0 configurations (#910), naturally produced deeply negative attention scores in every test (#743), and a mismatch budget for FP8 gradient compares at rounding boundaries (#947). The legacy test_mhas.py was removed as superseded (#986).
• Declared Python APIs are inspected in CI via an API index (#972), moved to test/api_index (#997) and then under test/python (#998), refreshed from the development CI wheel (#1008).
• Benchmarks: SDPA FLOP counts for causal masks computed without allocating masks (#347), a CUDA-event fallback when CUPTI cannot attach (#824), cudnn_oss benchmarked for kimi_k3/deepseek_v4 and FA2 on Ampere (#852), the unused Amax_S output dropped from the FP8 graphs (#847), artifacts refreshed against backend 9.26.0.39 (#771), a benchmark config-selection fix (#870), and README image links fixed (#772).
• CI: PR checks skipped on draft PRs (#770), an sm103 FROST SDPA failure fixed (#753), a conv ReLU execution-plan creation test fixed (#949), torch.matmul used instead of torch.einsum in test_mhas_v2 references to avoid SM107 worker crashes in cuBLAS (#952), causal-conv contract tests kept compatible with CuTe DSL 4.6.2 (#936), FP8 dprob reassociation drift allowed on Rubin (#939), zero-length sequences waived below cuDNN 9.25 (#944), a non-functional CUDA IMA guard removed (#791), guardword violations fixed (#929), and sensitive flags removed (#907, restored selectively in #923).

Bug Fixes 🐛

• A d256 FP8 + MXFP8 masked-softmax TMEM race (#985, issue #981) — the aliased P store raced the other warpgroup's S load. mhas_v2 forward now draws d=256 for FP8/MXFP8.
• d=192 THD decode int32 overflow (#984, issue #980) — the THD view's batch stride is now bounded, and the mhas_v2 decode sweeps draw THD.
• SM120 THD forward NaN poisoning from the K/V capacity tail (#884), and the f16 THD K/V descriptor extents clamped to the packed total with ragged capacity tails NaN-poisoned in test_mhas_v2 (#646, issue #624).
• The SDPA backward's stage 3 is compiled for the execution device (#937).
• The d256 FP8/MXFP8 SM100 kernels' import of the renamed THD helpers was fixed (#888, #893).
• Empty d192 quantized O TMEM loads are guarded (#757).
• SM120 prefill setmaxnreg budgets were rebalanced to fix d=256 register spills (#840); the cutlass prefix is kept for the SM120 backward (#903).
• A workaround was added for CUTLASS DSL 4.8's create_tensor_map_tiled (#839), and another to avoid an overflow (#872).
• The FA2 no-window call was fixed for flash-attn 2.7+ (#873).
• Nightly CI failures fixed (#789).

Removals 🧹

• The C++ NVRTC SDPA prefill OSS engine is removed (#792).
• The experimental torch SDPA op is replaced by cudnn::sdpa_fwd / cudnn::sdpa_bwd (#780).
• The legacy test_mhas.py is removed, superseded by test_mhas_v2 and sdpa/suites (#986).
• HSTU cross-tensor storage overlap checks are removed (#1012).

Acknowledgements 🙏

Thanks to everyone who contributed to this release:

@Adnios, @adshen, @Anerudhan, @Aneureka, @brandonfzhang, @Butterfingrz, @cb521, @egilliam-nv, @fallintoplace, @hwanseoc, @icavan, @JacoCheung, @jhjpark, @jiayus-nvidia, @Jie-Fang, @ksivaman, @liujane-dev, @miyanyan, @pmdavies-nv, @RomanAnders90, @SolitaryThinker, @SuperGoodGame, @swalters22, @thlurte, @thynics, @tp5uiuc, @vedaanta, @wanyingw, @YangXu1990uiuc, @yanqinz2, @yanzhuo607, @yihuawei, @yujincheng08, @yuweih205, and zhaorunchu@gmail.com.