cuDNN Frontend v1.30.0 Release Notes

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

New: the FlashInfer serving path runs on FROST 🚀 🚀

v1.29.0 shipped paged KV caches on the FROST SDPA engine and noted that FlashInfer already builds exactly the graph the engine wants. v1.30.0 is the release that makes that claim hold end to end: the graphs FlashInfer actually builds are re-declared in-tree as conformance suites, every place the engine read the buffer instead of the declaration is fixed, and the remaining declines on FlashInfer's serving shapes are gone.

The graph declaration is the operand contract (#1028). The cuDNN backend reads a variant pack as pointers, so callers bind buffers whose own shape need not match the declaration — a 2-D matrix to a [1, m, k] tensor, a flat quantizer blob to a reordered scale tensor, a 0-d scalar to (1, 1, 1). Python engines that read the pack saw the buffer's shape instead, so the same execute() succeeded on a backend plan and refused on a FROST plan. _normalize now describes a slot from the declaration whenever the caller's geometry disagrees with it but covers its bytes, with a precise rule: a buffer carrying the declared extents keeps its own strides (the linear-attention engines serve strided inputs that way), and a buffer of other extents is re-described only when its slots form one dense run that covers the declared bytes — which a transposed view of a contiguous block, such as FlashInfer's column-major B, does. The declaration supplies the dtype too (#1035): FlashInfer binds packed fp4 data as uint8 and the E4M3 scale blob as uint8, so a buffer whose storage slots are as wide as the declaration's is read as the declared dtype, while a narrower or wider buffer keeps its own and is never re-described.

The SDPA THD gate is over (H, S, D), not BSHD (#1028). Under ragged offsets the batch stride of Q/K/V/O is never read — every sequence base comes from the offset table and the lowering binds the batch axis at extent 1 — yet the engine gate and the DSL adapter required BSHD-physical order over all four axes. FlashInfer declares the batch stride equal to the token stride (h * d), so its ragged prefill was declined at every b > 1. packed_layout_ok now checks head dim innermost, then heads, then tokens, and the same predicate is used by the mismatch gate and by both the SM100 and SM120 adapters.

FlashInfer's padded (b, s_max, h) Stats is served under THD (#1036). A THD graph whose Stats tensor has no ragged offsets is the per-batch padded form — FlashInfer's return_lse buffer. The packed path wrote token rows contiguously, so at b > 1 the form was declined and at b == 1 it coincided by luck with the rows past a sequence's length left unwritten, which is what failed 52 of FlashInfer's token_indptr tests once an earlier test had dirtied the allocator. The 20 SM100/SM107 THD templates now take lse_padded_rows and select the per-batch form on the static fake, and the adapter seeds the buffer with -inf on the launch stream, matching the backend's contract. Nothing is added to the kernel ABI.

THD templates compile with dynamic batch and head extents (#1039). A compile() pinned b, qh, kh and every stride in its fakes, so FlashInfer's serving shapes minted one 2.6 s kernel per (b, qh, kh, strides) — 126 forward kernels in FlashInfer's cuDNN attention tests. The kernels never needed it: _host reads B/QH/KH from problem_size at run time. compile(dynamic_bhk=True) rebinds them to cute.sym_int() right after the cache key is taken, giving one kernel per layout class.

Dense padded-Q trim on every kernel (#1037). FlashInfer's dense batched prefill hands per-batch Q lengths with every padded graph. The SM100 fp8/mxfp8 d128 and (192, 128) flavors previously either ran a host read (seq_len_q.min().item(), which breaks CUDA-graph capture) or were declined; all four remaining kernels now carry the same trim as prefill_d256_fp8, and the capability row and the execute-time host read retire.

Split-KV is also kept on paged decode when the declared KV max is not a 128-multiple (#1092) — FlashInfer passes its true max verbatim, and a paged graph is padded by construction, so the clause that declined it never applied.

New: persistent compiled-plan cache 🚀 🚀

cute.compile runs the whole DSL backend once per process per distinct kernel — 0.7 s for a FROST GEMM, 2.6 s for an SDPA prefill on SM100 — and the DSL's own file cache stores only MLIR bytecode. cudnn.frost.compiled_cache (#1031) keeps the exported tvm-ffi object of every kernel compiled with --enable-tvm-ffi and reloads it in milliseconds, across processes.

• compile_cached(fn, *args, cache_key=, symbol=, **kwargs) is a drop-in for cute.compile; hit and miss hand back the same reloaded artifact.
• Identity is the content hash of a manifest — frontend, cutlass-dsl, tvm-ffi, CUDA driver, device name/CC/SM count/L2. An entry is reused only under its own embedded key, the object is written before its record (temp file + os.replace), and anything doubtful is a miss. An environment with any unknown manifest field neither loads nor persists.
• The eight GEMM kernel templates and the 49 cute.compile sites across the 36 SDPA forward/backward template files route through it. GEMM suites: 121 tests, 44 s cold → 12 s warm, and a second process reloads and computes a bit-identical result.
• prune() (#1038) retires whole dead environment directories — dead schema roots first, then oldest first, never the process's own — under CUDNN_FRONTEND_COMPILED_CACHE_MAX_BYTES (4 GiB default, 0 disables), run once per process after its first export. Prune never follows symlinks, never leaves the root, and removes only directories this cache made.

New: one knob vocabulary for backend and Python plans ✨✨

A Python (FROST) plan is now described, named and replayed exactly like a backend plan: (engine_id, {cudnn.knob_type: int}) (#1026).

• knobs.h freezes the 33 backend-mirrored KnobType_t values and adds a frontend-only band from FRONTEND_KNOB_TYPE_BASE = 1000 (SCHED_POLICY, PACK_GQA, SPLIT_KV) that convert_to_backend_knob_type refuses. Both bands are append-only, pinned by static_assert.
• Knobs are performance-only. The numerics-changing softmax accumulator precision leaves SdpaFwdKnobs and becomes the sdpa() / sdpa_fp8() op attribute softmax_precision, which is never forwarded to C++ and makes the node backend-unlowerable when set.
• get_engine_and_knobs_at_index always returns a dict, create_execution_plan accepts the public dict, and get_plan_name_at_index renders engine[TILE_M=128, ...].
• The frost_gemm family gains the facts + heuristics hooks so its plan enters the ranked list with the TileConfig it will build; GemmKnobs ↔ TileConfig is lossless through the canonical config name, so a recorded (engine_id, knobs) replays the exact kernel. A record naming no canonical config is declined, never snapped.
• The SDPA-backward family gains a recommend() hook, so its plans list the tiles the lowering will pick rather than {}.
• pygraph.serialize() now applies key()'s rule and raises cudnnGraphNotSupportedError on a backend-unlowerable graph before any lowering — previously a SET softmax_precision serialized as the f32 pipeline and would deserialize and run with different numerics.

Plan ranking improved alongside it: FROST SDPA plans are ranked by workload against measured shards (#1152), SM100 heuristics account for candidate launch geometry (#1158), and the FROST GEMM heuristic takes the 64-byte-K MMA and the 256-tall CTA pair on the parts that issue them (#1177, #1150).

New: decode tiles — SDPA decode stops running the prefill pipeline 🚀 🚀

The SM100 f16/bf16 row served paged decode with the prefill pipeline: 512 Q rows per cga2 cluster with 1..G of them live at S_q = 1, so every KV tile paid full dead-row BMM1/BMM2 and two softmax warpgroups of exp on zeros. FlashInfer's Qwen3-235B decode shape (b=32, H=64/4, d128, S_kv=4096, page 16, bf16) measured 118 µs against 44 µs for the cuDNN backend's decode engine.

• d128 decode tile (#1094) — sm100/decode_d128_f16.py, TILES_Q=1 at cga1, one softmax warpgroup, Q/O SMEM aliased so the K/V ring is three stages deep (224 KiB), and two S/P TMEM slots alternating per KV tile so BMM1(i+1) overlaps softmax(i). Same contract and ABI as the prefill kernel, paged pools included.
• d256 decode tile (#1109) — sm100/decode_d256_f16.py, swap-AB, selected for decode-shaped graphs (S_q × PackGQA group ≤ 32 rows: Qwen3.5 32/2 at S_q ≤ 2, Qwen3-Next 16/2 at S_q ≤ 4, MHA at S_q ≤ 32). Same engine row and decline discipline as the prefill d256 tile.
• Partial PackGQA (#1104) — a GQA group G = H_q / H_kv that does not divide TILE_M could not pack at all, so a 96/8 (G = 12) d128 paged decode ran unpacked with one live row per 512-row cluster and the KV head re-read 12×. The d128/d256 f16 kernels now pack p = gcd(G, TILE_M): 96/8 packs 4 heads per token row-group, 48/8 packs 2. A group sharing no factor with the tile still cannot pack, and a pinned PACK_GQA=1 stays declined rather than silently running unpacked.
• Attention sinks at S_q == 1 (#1095), including paged KV and sliding window — the python-native validator's blanket "decode only mode not supported with sink_token" rule is dropped, since whether a sink at S_q = 1 is served is an engine's support-surface answer at planning time.
• Paged KV on the d192×128 f16/bf16 flavor (#1096), with paged head dims gated on the selected flavor. The K pool is 192 wide and the V pool 128 wide — the Kimi-Linear / DSv3-style non-absorbed MLA decode shape.

New: prepared launches — the execute fast path 🚀 🚀

The f16/bf16 SM100/SM107 forward kernels move onto an explicit pointer/int host entry, so every extent and stride is a runtime argument and the compile key is layout-only (#1119). lower_dsl_prefill resolves at plan time what the graph fixed — operand ids, IR layouts, which feature operands the facts demand, the quantized-operand id table, a raw-stream → CUstream cache — and _execute_resolved then does dictionary lookups and one api.execute(**kwargs).

• Execute-time overrides apply in one native crossing (VariantPackNative.override_many, #1132), with the graph keeping a uid → slot map next to its cached DeclaredLayout. On a FlashInfer-shaped THD graph, a bounded batch override goes 73–75 µs → ~35 µs; the backend plan pays the same loop and benefits alike. Dense launches join THD on the prepared positional entry in the same change.
• The declaration rule itself runs natively, once per pack (#1028) — as first written it cost +7.5 µs on both the backend and FROST paths and +10 µs more for a 2-D binding; it is now back to baseline, and a 2-D binding costs what a declared one does.
• Int64 stride leaves (#1146) fix a 2^27 TMA batch-stride overflow introduced by #1119: the DSL typed the (batch, seq, head) stride elements as dynamic Int32 from the compile() placeholders, so a 16-bit operand with S·H·D == 2^27 elements per batch encoded a negative batch stride (launch aborts) and 2^28..2^30 encoded 0 (every batch silently aliases batch 0). The stride placeholders are now cutlass.Int64 and _bshd coerces its leaves, so every TMA descriptor is 64-bit whatever the caller bound.
• Further overhead work in #1151: memoized THD launch geometry, a bounded LRU for dense layout validation, prepared buffer facts packed in the native variant pack, prepared split-KV launches, the d128 decode pointer-ABI migration, and FE resource cleanup made safe during CUDA-graph capture.
• Dense forward reads strided Q/K/V/O zero-copy (#1102). The dense path normalised any non-BSHD-compact operand with a transpose plus .contiguous() before a descriptor was ever built — a full gather of Q on every execute, and exactly what a caller slicing a fused QKV projection hits. Nothing in the kernels needed it: they address Q/K/V only through TMA coordinates. Twelve kernels across SM100 and SM107 now decide per operand at compile time.

New: gated attention block 🚀 🚀

The SM107 d256 epilogue gains a fused gate, O *= sigmoid(G), in the production kernels and as an op graph, with a composable gated attention block on top (#1102). It grew over the release into a full quantized block:

• Optional QK-RMSNorm, MXFP8 unfused and fully fused, and an MXFP8 SDPA gate (#1110). sdpa_mxfp8 also learns to infer O / Stats / Amax_O dims the way sdpa_fp8 does, so a caller that never asks for Amax_O no longer fails Tensor.validate().
• MXFP4 projection weights and an optional fp4 gated output (#1117) — MxQuantSpec.w_qkvg_dtype puts MXFP4 weights on the mixed block-scale GEMM row, and MxQuantSpec.o_fp4 = NVFP4 | MXFP4 block-quantizes the gated O to feed an fp4 × fp4 block-scale out projection. New tile_dsl fp4 primitives (fp32_to_fp4_pack via cvt.rn.satfinite.e2m1x2, e4m3_scale_from_amax, e8m0_from_amax) and a quantize_fp4 kernel writing per-token e2m1 codes directly in the out-projection GEMM's F8_128x4 order. proj_gemm takes a per-operand dtype, block size 16|32 and E4M3|E8M0 scale dtype through one pairing table derived from the FROST block-scale catalog.

New: block-scaled O epilogues on the quantized SDPA forward 🚀 🚀

The per-tensor FP8 forward can now emit O as FP4_E2M1 (two per byte, one E4M3 scale per 16 d elements) or as FP8_E4M3 with one UE8M0 scale per 32 d elements, writing the scale factors to a new optional sdpa_fp8 output sf_o in F8_128x4 atom order (#1088). The epilogue reuses the row-owning correction warps on the d128 SM100/SM107 kernels and a quad butterfly on the SM120 kernel; scale_o doubles as the FP4 global scale.

The MXFP8-input forward gets the same contract on the SM100 d128 kernels (#1180), including a python-only scale_o input — required for an FP4 O, since the E4M3 block scale alone cannot span its range — and rejected without sf_o.

New: DeepSeek-V4.1 building blocks 🚀 🚀

A family of prepared FROST primitives for DSv4.1:

• Saved-state Engram gate, forward and backward (#1113).
• Tail RoPE and group32 FP4/FP8 QDQ (#1114), extended with group16 E4M3-scale FP4 (#1115), and a prepared BF16 tail RoPE for prefill (#1116).
• Fused native-layout vision RoPE backward with QKV gradient packing (#1122).
• Prepared mHC projection / RMS backward (#1179), with the public exports indexed and the published API navigation updated.

New: FROST SM100 convolution forward 🚀 🚀

[FROST] Add sm100 conv forward (#961) brings the first convolution kernel to the FROST engine family on Blackwell.

New: Hopper (SM90) Kimi Delta Attention 🚀 🚀

There was no KDA path on Hopper at all: the FROST kernels are Blackwell-only by construction (42 tcgen05 and 84 tmem references in kda_prefill_f16.py alone), and the only other backend, cuTile, needs the cuda.tile runtime — so on an H100 all three linear-attention ops raised cudnnGraphNotSupportedError. Relaxing the arch gate cannot work, because Hopper has no Tensor Memory and no tcgen05 MMA: SM90 needs a different schedule, not a port.

• A CuTe DSL prefill path (#1017) built on warpgroup (wgmma) against shared memory, with a chunk-parallel PREP pass feeding a sequential SCAN over the [128, 128] state, and a mid-chunk anchor that keeps every exp argument inside ±40.
• A second engine backed by a fused CUDA kernel, forward and backward, compiled under NVRTC, plus a CuTeDSL kernel swap and default-engine heuristics (#1061). The descriptor rename dkg → d_kg followed in #1103.

Sparse attention: DSA and BSA 🚀 🚀

DeepSeek Sparse Attention (DSA).

• SM100 H96 backward specialized as H64 + H32, always routed to the composite (#1011).
• SM100 H64/D576 backward scheduling and dKV writeback optimized, with compact FP32 dKV rows kept after GB200 measurements showed no benefit from a 640-element stride (#1014).
• SM100 2-CTA backward for BF16 H128/D576 (#1048) — a (2,1,1)-cluster CuTe DSL kernel where the two CTAs split the 128 heads on cta_group::2 tensor cores, P and dS exchanged with shared::cluster bulk copies, and slot validity applied before every score and gradient operation so ignored top-k slots cannot contaminate gradients.
• SM100 indexer score forward and recompute optimized (#1147).
• CPU overhead reduced in the DSA and BSA launch paths (#1073), and a config-driven DeepSeek-V4 sparse attention benchmark suite added (#1067).

Block Sparse Attention (BSA).

• Native SM120 blk128 forward with an FA4-style path, four-fold unrolling and a paired performance gate (#1070).
• Optimized SM90 blk128 forward and backward (#1153) — native Q128 forward with independent K/V pipelines, overlapping attention compute and split-KV support; Q64×KV128 backward with two compute warpgroups, direct TMA dK/dV stores and independently pipelined LSE/delta prefetch. Validated under Compute Sanitizer memcheck, racecheck and synccheck.
• SM100 blk128 JAX BSA with an explicit backward and custom VJP (#976), followed by two errata passes that fix the bridge regressions and expose the APIs under cudnn.jax and cudnn.torch (#1081, #1111).

Rubin (SM107) 🚀 🚀

• MXFP8 prefill closes the gap to FP8 (#1059). The d128 and d192×128 SM107 MXFP8 kernels ran ~10 % behind their per-tensor FP8 siblings, and the whole deficit was the softmax warps' register-file row-sum — 63 packed adds per row per KV iteration, which makes the compiler sink the alpha/stats and first P-chunk publishes to the end of the exp burst so BMM2 waits on softmax every iteration. The row-sum now rides the tensor core: one N=16 MMA of the e4m3 P operand against an all-ones SMEM tile. Zero-copy SF bind, scheduler claims and causal ranking land with it.
• +5..17 % on fp8/mxfp8 prefill from the Amax_O fold (#1129). cute.math.max lowers to arith.maxnumf, which the DSL → NVVM path emits as a COMPARE + SELECT pair, costing three ALU ops per O element in the epilogue's critical tail — the d512 fp8 kernel carried 1548 FSETP + 1541 FSEL per tile against the C++ reference's 6. tile_dsl.pointwise.fmax_f32 emits PTX max.f32, which ptxas fuses into FMNMX/FMNMX3 with |x| folded into the operand modifier: one instruction per element, NaN semantics unchanged.
• One predicated scheduler credit arrive per warp (#1169). The previous spelling lowered to LDC + BRX jump tables + BSSY/BSYNC/BREAK reconverges and one arrive per arm — 51–56 instructions per warp per work item on the sm107 d512 fp8 prefill — and is now a single predicated mbarrier.arrive. A hint-less ring-wait spin becomes a per-kernel opt-in on 7 Rubin prefill flavors, with d512 fp8/mxfp8 kernel hoists.
• 512 CTA M size enabled (#1054), and the persistent THD scheduler consumed in the last two FP8 kernels (#892).
• MXFP8 MoE expert-parallel training (#750) — a new MoeEp Python API with validated forward and backward contracts and lazy optional-dependency loading, vendored MegaMoE CuTeDSL communication/workspace/scheduling primitives, Rubin forward GLU and backward dGLU training kernels, and a runtime resource layer managing NVSHMEM lifecycle, symmetric workspaces and capability checks behind a lazy backend seam.
• MXFP8 clamped GeGLU forward and backward (#1144), passing activation alpha and clamp bounds through the Rubin grouped GLU APIs and correcting non-unit GEMM alpha scaling in the GeGLU backward.

SDPA: more kernels and numerics 🚀 🚀

• A specialized SM120 D512 FP8 prefill kernel (#1120).
• PV_BF16 for the MXFP8 (QK) SDPA kernel (#983), avoiding causality leakage, with a D192 hybrid benchmark and FROST kernel-time profiling.
• SM80 THD / ragged backward through the graph API (#950). graph_analyzer.thd_stats_packing becomes the one classifier for packed Stats — token_major, head_major or None — replacing the backward probe's, the lowering's and the two forward adapter sites' own copies, and the SM80 backward kernel reads packed Stats in both packings with an arch-agnostic lengths → cu_seqlens launch.
• SM100 exp2 on the FMA pipe (#1178). A softmax KV iteration on a 128-wide S tile issues 128 ex2.approx per row, making MUFU the softmax warps' longest pipe while FP32 has slack. exp2_emul_pair evaluates two exp2 in 6 packed FP32/INT instructions with no MUFU — the same split the cuDNN backend kernel uses — and exp2_mixed routes a compile-time subset of a vector's pairs through it. Applied to d128 MXFP8, d128 FP8 and d192×128 bf16 prefill at cc 10.0.

New: stats_use_log2 — LSE in base 2 ✨✨

Flash-attention-style consumers (FlashInfer, FA2/FA3, TRT-LLM) define the LSE as max + log2(sum_exp); cuDNN returns max + ln(sum_exp), so callers convert with a separate elementwise pass — and a doubled conversion already shipped a silent 1.44×-wrong LSE with a bit-exact O (flashinfer-ai/flashinfer#4663). The convention is now owned at the graph level: SDPA_attributes::set_stats_use_log2 / graph.sdpa(stats_use_log2=...) (#931). The FROST forward engines serve it natively — every prefill kernel already keeps the softmax in the log2 domain, so each gets one const_expr-guarded lse *= log2(e).

Follow-ups fixed the FP8 Stats log-base binding, protected the SM120 split partials and covered native ragged decode GQA Stats writes (#1082). The attribute's final home is the unified softmax node (CUDNN_ATTR_OPERATION_SOFTMAX_STATS_LOG2), gated on cuDNN 9.27.0 (#1127); graph.sdpa(stats_use_log2=...) is unchanged.

Linear attention ✨✨

• state_indices for paged recurrent states (#1002). Serving stacks (vLLM, SGLang) keep the linear-attention recurrent state in a paged pool and hand the kernel one row id per sequence rather than a compact buffer in sequence order — FlashInfer's chunk_gated_delta_rule models this as state_indices. cuDNN had no equivalent, so a caller holding a pool had to gather the active rows in and scatter them back out around every call, costing more than the kernel saves. An optional mStateIdx is now threaded through the SM100 GDN prefill kernel; absent it the state row is the sequence, exactly as before.
• Context-parallel samples and performance improvements (#999), and P = #SMs/(B·H) = 2 tuning (#1164).
• Q normalization skipped for the FROST LA summary node, with everything directed to the fused path (#1105).
• Fixes: backward tail work assigned by CTA ID rather than SM placement, which is not unique under concurrent kernels (#1033, fixes #1032); KDA summary factor-warp parity preserved across work items (#1064); numerical issues in the Neumann inverse for KDA/GDN-2 (#1124); and a Hopper handle that must never be wrapped in torch.cuda.ExternalStream when it is the default stream (#1165).

GEMM and MoE ✨✨

• Canonical (natural row-major) layouts and flat SF buffers accepted by the contiguous grouped GEMM SwiGLU/dSwiGLU wrappers and API classes (#796) — A (sum_m, k), B (l, n, k) C-contiguous, dense C-contiguous SFA/SFB buffers flat or in physical atom shape. Canonical SF buffers compile as flat 1-D pointers, since the kernels already rebuild the MMA-tiled SF layouts from the GEMM shapes. Pre-permuted kernel-facing inputs keep working unchanged.
• Dual Torch/JAX SwiGLU wrappers and namespace APIs (#1079), sharing one entry point and dispatching to either framework.
• FROST GEMM code refactor plus MoE matmul (#1068), MoE swap-AB (#1090), and reduced CPU overhead in MXFP8 MoE training by caching block-scaled GLU/dGLU and wgrad wrapper setup (#924).

Graph API, Python API, and other operations ✨✨

• Torch custom ops for LayerNorm and RMSNorm (#1128), with lazy exports.
• Optional max-logit output for Flex Attention (#1131) — per-query-head maximum scaled logits across the SM90 and SM100 forward kernels, exposed as return_max_logit with preallocated outputs, stream ordering, compile-cache specialization and non-differentiable autograd results. A Q-stage-2 empty-tile bug that retained NaNs from reused TMEM is fixed by writing literal zeros.
• AGENTS Rule 8 (#1167): the graph API owns no device memory and never blocks the host, at build as well as at execute. The rule names the ownership and blocking bans for both phases, documents that a dead ABI slot may be 0, and names the CAKE decline-under-capture shape as the only legal exception. #1168 applies it: one shared cudnn._torch_stream helper replacing eighteen local copies across sdpa, linear_attention, conv1d, flex, hstu, grouped GEMM, CSA, DSA and QAT; no plan-owned device memory in the GEMM/FROST MoE plans; dead ABI slots set to 0; and no host sync in the SWA execute path.

Documentation 📖

• A README for the Fern documentation tree (#1060), with the intro tightened for the external repo (#1062).
• A stale benchmark doc reference updated and kernel comments reworded (#1041).
• The convolution mapping docstring clarified to avoid a guardword match (#1076).
• Documentation and benchmark images converted to lossless WebP (#1078).

Tooling, CI, Benchmarks, and Tests 📊

• Recurring CI reds turned into tracked xfails or fixes (#1162), and OSS CuTeDSL test warnings cleaned up (#1161).
• JAX tests skipped where XLA cannot compile for the GPU, with worker deaths made legible (#1055).
• Gated-attention test modules gated on the FROST DSL requirement so a too-old nvidia-cutlass-dsl reads as a version skip rather than an import failure (#1121).
• Paged SDPA properties (#1063) and SM80 backward THD extent properties (#1077) recorded in the API index.
• SDPA fuzzing over head-axis stride gaps in the ragged sweeps (#960); Rubin exact-LSE tests moved to an fp64 reference (#1075); the paged SM100 sync-debug mode armed inside the CUDA-graph capture rather than around it (#1053).
• Benchmarking artifacts refreshed against backend 9.27.0.21 (#1176).
• The block-mask sample skipped on SM10x with cuDNN < 9.26 (#1189), plus general CI fixes (#1190, #1118).

Bug Fixes 🐛

• HSTU LMSD backward coordinate type join (#1057) — reusing the first predicated loop's row_tile_coord for the dX output loop fails with TYPE_UNSTABLE_JOIN when Quack 0.6.4 rewrites mixed constexpr/runtime guards under CUTLASS DSL 4.8. The dX loop gets its own dx_tile_coord; coordinate, predicate and arithmetic are unchanged.
• TMA batch-stride overflow on the f16/bf16 pointer host ABI (#1146) — see prepared launches above.
• SM90 DSA backward dS scaling precision (#1046) and SM90 indexer dK staging synchronization (#1049).
• SM100 dense indexer dK staging synchronization (#1135) and SM100 H16/H32 DSA backward O*dO operand precision (#1134).
• SM90 singleton score layouts and causal masks in DSA (#1154).
• D256 MXFP8 backward synchronization fixed with dQ/dKdV optimized (#1040).
• Output alignment check fixed (#1074).
• A series of assorted fixes (#1118) and CI repairs (#1190).

Acknowledgements 🙏

Thanks to everyone who contributed to this release:

@Anerudhan, @Aneureka, @brandonfzhang, @C-TC, @Denny991, @egilliam-nv, @elfiegg, @hwanseoc, @icavan, @jhjpark, @jiayus-nvidia, @mdy666, @NVIDIA-JerryChen, @pmdavies-nv, @rmhaskarnvidia, @RomanAnders90, @thynics, @tiffany940107, @vedaanta, @XinboZhao, @YangXu1990uiuc, @yanqinz2, @yanzhuo607, @yihuawei, @yuweih205, @zach-ye0, @ZeYang1025, and @zhibinz-nv.