1.28.0 rc

[Anerudhan]

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 (Add an opt-in cuDNN FLA GatedMLP shim #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 (SM80 (A100) SDPA: FROST engines + cudnn.sdpa adapters #493), later ported to plan-time compilation with TemplateParams kernels and sym_int THD extents ([SDPA] SM80 fwd: TemplateParams kernels, plan-time compile, sym_int THD extents #689). Previously the manifest had no engine below SM100.
• SM120 — a forward per-tensor FP8 (E4M3) engine (Add SM120 per-tensor FP8 (e4m3) SDPA-forward engine #509) grown through several rounds of features and tuning (frost(sdpa): add more features like padded head dim, skv_tile=0, diagonal_band_right_bound to sm120 frost sdpa prefill kernel #531, frost(sdpa): optimize and enhance sm120 sdpa forward fp8 kernel #562, frost(sdpa): add more features to sm120 frost fp8 sdpa_fwd kernel #595), MLA d_qk=192 / d_v=128 support (frost(sdpa): add mixed head-dim support for SM120 frost sdpa_fwd engine #507), and a backward engine (Add SM120 FROST SDPA backward engine (sdpa_bwd_sm120) #486) extended to all head sizes ≤256 including d192/d256, non-compact layouts and deterministic dQ (frost(sdpa): SM120 backward — all head sizes (≤256), d192/d256, non-compact layouts, deterministic dQ #533), sliding-window attention (Add sliding-window attention support to the SM120 FROST SDPA backward engine #505), GQA/MQA, padding masks, right-band widening and sink-token gradients (frost(sdpa): GQA/MQA, padding mask, right-band widening, and sink-token gradients for the SM120 f16 backward #557), deterministic 2-kernel mode and dBias (add det 2k and dbias support for sm120 sdpa bwd #707), and native service of declared strided layouts (frost(sdpa): serve declared layouts natively in the SM120 backward — strided stats/io, TMA zero-fill head-dim envelope #666).
• SM107 (Rubin) — per-tensor FP8 prefill via a sibling kernel (sdpa: enable per-tensor FP8 prefill on SM107 (Rubin) via a sibling kernel #576), with the has_lse specialization and a static SMEM guard (sdpa fp8 sm107: port the has_lse specialization; add a static SMEM guard #579), a fused LDTM row-max and row-sum-in-MMA epilogue (sdpa fp8 sm107: fused LDTM row-max + row-sum-in-MMA #580), and the softmax_precision knob axis lit up with an F16x2 exponent on the d128 sibling (sdpa fp8: light up the softmax_precision knob axis — f16x2 exponent on the d128 SM107 sibling #651).
• MLA backward — sdpa_bwd gains MLA support (feat(frost): add mla support for sdpa bwd #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 (feat(sdpa): accept max_total_seq_len_q/kv on the forward SDPA 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 (sdpa: require FP32 Stats and default an unset Stats dtype to 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 (fix: reject non-BHSD softmax-stats strides on cuDNN < 9.26; randomize in tests (NVBug 6057616) #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 (fix(sdpa): move the pre-9.26 Stats packed-BHSD check to post_validate_node #642).
• Capabilities.bottom_right_padded_seq_q was retiled (frost(sdpa): retire Capabilities.bottom_right_padded_seq_q #683).

Serialization and plan management

• Handle-less plan deserialization ( feat: handle-less plan deserialization via DeviceProperties (cuDNN 9.8+) #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 (fix : serialize Tensor_attributes::alignment #564).
• CUDNN_KNOB_TYPE_TILE_CGA is mapped (Map CUDNN_KNOB_TYPE_TILE_CGA so knob queries and explicit plans round-trip #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 (Rank plans in one place: delete propose_plans, Router, and heuristics_sort #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 (Normalize the variant pack once, into a C type that is also the DLPack producer #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 (Make a family a kind of graph, and stop paying for the DSL to decline #502).

Operations

• FFT causal conv1d bindings are back (Reintroduce FFT causal conv1d frontend bindings #542). Reintroduces the support merged in Add FFT causal conv1d frontend bindings #437 and reverted in Revert "Add FFT causal conv1d frontend bindings (#437)" #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 (Skip SM107 causal conv1d samples before cuDNN 9.26 #632).

Open-Source Kernels 🚀 🚀

GEMM and MoE fusions

• cudnn.gemm.ops.swiglu_mlp (Add a dense BF16 SwiGLU MLP autograd op with fused forward and dSwiGLU backward #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 (Add block-scaled grouped GEMM + SwiGLU + RHT + NVFP4 quantization fusion for Rubin #637), built on new E5M3 scale-factor support for fused GEMM on Rubin (Add support for E5m3 fused GEMM on Rubin  #545); the GLU+RHT+quant kernel now emits column-wise RHT in ragged tensor layout (GGEMM+GLU+RHT+quant kernel now outputs column-wise RHT in ragged tensor layout #693); FROST GEMM gained Rubin kernel pipelines (frost(gemm): Add Rubin kernel pipelines #593); and single-group runtime offsets were added to Rubin grouped dGLU (Add single-group runtime offsets to Rubin grouped dGLU #590).
• Grouped SiTU-GLU activation ([CuTeDSL] Add grouped SiTU-GLU activation #645), with the SiTU-GLU Hadamard API contracts fixed ([CuTeDSL] Fix SiTU-GLU Hadamard API contracts #670).
• Deterministic dprob for grouped GEMM dsrelu (feat(determinism): grouped gemm dsrelu deterministic dprob #521) and FP32 row-scaled FP4 grouped GEMM (Support FP32 output and dynamic M in row-scaled FP4 grouped GEMM #461).
• Host overhead cut roughly 3–4× on the CuTeDSL grouped-GEMM wrappers (reduce cuteDSL grouped gemm CPU latency #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 (Make the public gemm path work for the flavors it claims #558), with non-MoE, workspace-free flavors lowered off one build-time recipe (Lower non-MoE, workspace-free gemm flavors off one build-time recipe #559).
• Kernel-level throughput work: mixed CGA for the FROST GEMM kernels (Mixed CGA #617), multiple tcgen05 instructions in the M dimension (Multi tcgen05 instructions in M dim #566), multi-UTCMMA optimization (Multi utcmma opt #511), the LPT_L2 scheduler (Add LPT_L2 scheduler #585), TMA-STG with multi-GEMM (Tmastg with multi gemm #721), M-major output refactorization (M major output refactorization #735), and Rubin MoE GEMM optimization passes (Rubin perf investigate #667, Further rubin optimization #697, Optimize rubin moe gemm #700).
• SM120 matmul support (Add SM120 Matmul Support #719).
• The NVRTC dependency for occupancy queries was removed (remove-occupancy-nvrtc-dependency #494), and the elect_one compilation hint optimized (optimze-elect-one-compilation-hint #504).

DSA (DeepSeek Sparse Attention)

• Opt-in SM100 sparse indexer backward v2 (dsa(indexer_backward): opt-in SM100 sparse backward v2 — 1.16-1.92x faster (fp32-accurate d_index_k at no extra cost) #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 (Improve performance of SM100 DSA backward kernel #684), and an SM100 H16 backward specialization added for H=16, d_qk=576, d_v=512 BF16 (Add SM100 H16 DSA backward specialization #664).
• indexer_backward now validates the output and plan signature before kernel 1 on the default SM100/SM90 backends (dsa(indexer_backward): validate output & plan signature before the score-grad precompute (default SM100/SM90) #572), and range_constexpr was restored in eight kernel_gemm epilogue loops (dsa(indexer_backward): restore range_constexpr in 8 kernel_gemm epilogue loops #549).

Block-sparse attention (BSA)

• Sage FP8 forward for Blackwell (BSA: add Sage FP8 forward support 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 (Enable CSA compressor on SM100+ GPUs #641), with the validated ratio envelope extended to ratio=2 (CSA compressor: extend the validated ratio envelope 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 CSA compressor: review-response fixups for the ratio=128 kernels (follow-up to #427) #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 ([1.28.0-rc] Gate the packed-FP4 wgrad layout workaround on cutlass-dsl < 4.8 #764).

Tooling, CI, and Build ✨✨

• Style checks moved to GitHub Actions (ci: enforce pre-commit style check via GitHub Actions #741) — pre-commit now runs as a GitHub Actions job.
• Merge requirements are enforced (ci: require Milestone and Projects fields on PRs before merge #736, ci: report a clear error when the Projects lookup fails in merge-requirements #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 (Add stale-PR automation: mark stale at 3 weeks, close at 6 weeks of inactivity #714).
• setup.py defaults to a parallel extension build (build: default to a parallel extension build in setup.py #565), and a -Werror unused-parameter build break in init_gnn_submodule was fixed (python: fix -Werror unused-parameter build break in init_gnn_submodule (pre-9.26 headers) #728).
• The large-tensor fuzzer's plan fallback and failure output were improved (Improve large tensor fuzzer plan fallback and failure output #583), on top of a new large-tensor convolution fuzzer (Add large-tensor convolution fuzzer #401).

Samples, Benchmarks, and Tests 📊

• Attention inference benchmark suite (benchmark: attention inference suite (context/generation, cudnn vs cudnn_oss) #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 (Add reproducible Qwen3.8 and Qwen-Image 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 (Fix packed-QKV views in the cuDNN FLA GDN shim #685), the cudnn.fla GatedMLP shim (Add an opt-in cuDNN FLA GatedMLP shim #686), and the backend-only public d256 SDPA after SDPA: drop the legacy standalone d=256 fwd/bwd stacks; port SM80 forward to the SdpaFwdDsl adapter path #682.
• The attention benchmarks were restructured (Restructure attention benchmarks and update benchmarking artifacts (backend 9.26.0.17, frontend b21cd3f) #657) and the benchmark scripts unified (Benchmark scripts unification #655); the SDPA benchmark gained the cudnn_oss (FROST) backend with unified sustained-clock SOL and peak lines (benchmark/sdpa: add cudnn_oss (FROST) backend, sustained-clock SOL + peak lines; refresh gb200/gb300 and add RTX PRO 6000 results #597), FA4 auto split-KV enabled by default (benchmark(sdpa): enable FA4 auto split-KV by default (num_splits=0) #607), a Qwen3-VL vision-encoder (ViT) config with GB300 results (benchmark: add Qwen3-VL vision-encoder (ViT) SDPA config + GB300 results #598, benchmark(sdpa): qwen3vl_vit writes to results/ like every other config #629), SM-clock sampling on the GPU the benchmark actually runs on (benchmark: fix peak-MMA/SOL clock sampling under CUDA_VISIBLE_DEVICES (peak-MMA line collapse on multi-GPU nodes) #699), and an Ampere (sm80) row in the peak-MMA table (benchmark: add the Ampere (sm80) row to the peak-MMA table #715). GDN benchmarking was added in Add GDN benchmarking #501.
• Benchmarking artifacts refreshed against backend 9.26.0.17 and 9.26.0.33 (Restructure attention benchmarks and update benchmarking artifacts (backend 9.26.0.17, frontend b21cd3f) #657, Update Benchmarking Artifacts - backend 9.26.0.33, frontend 73d8feb #665, Update Benchmarking Artifacts - backend 9.26.0.33, frontend 55b2773 #681, Update Benchmarking Artifacts - backend 9.26.0.33, frontend d53cafc #694, Update Benchmarking Artifacts - backend 9.26.0.33, frontend aded990 #711).
• Tests: ragged per-tensor token-stride gaps are fuzzed in test_mhas_v2 (test(sdpa): fuzz per-tensor ragged token-stride gaps in test_mhas_v2 #516) and sink tokens are drawn in the ragged backward suites (test_mhas_v2: draw sink tokens in the ragged bwd suites #630); DSA comparisons use only the effective top-k slice (test(dsa): compare only the effective top-k slice #672); render-only tests were removed (remove-render-only-tests #688); samples skip an unsupported case and drop an invalid cudaGraphDestroy (samples: skip unsupported test case, remove invalid cudaGraphDestroy #734); the 02_low_level_api notebook uses an FP32 Stats tensor (samples: use FP32 Stats tensor in 02_low_level_api notebook #701).
• CausalConv1d sample links and the backends architecture fence were fixed (Fix CausalConv1d sample links and backends architecture fence #518).

Bug Fixes 🐛

SDPA

• Ragged backward workspace under-allocation with head-major stats (Fix ragged SDPA backward workspace under-allocation for non-token-major stats layouts #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 (sdpa: fix SM100 frost zero-KV cluster deadlock + enable zero-length seq_len_kv sweeps #575).
• Fixed a latent THD O-descriptor row stride on the SM100 FP8 path (sdpa fp8 sm100: fix THD O-descriptor row stride (latent) #577).
• Fixed SM100 D192 E5M2 attention-sink accuracy (Fix 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 (fix: bind a CUDA context 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 First-class cudnn.Handle (create_handle returns an object owning {backend handle, device, stream}) #612 were added in python: add the ensure_current_context #612 imports but never defined #638.
• The oversized-SMEM ceiling is requested by ordinal when the binding cannot name the device (Fix device crash #635).
• DLPack is told the frontend is only reading, so probe() survives a graph capture (Tell DLPack we are only reading, so probe() survives a graph capture #506).

Kernels

• Three Rubin (sm_107) compile and launch failures fixed (Fix three Rubin (sm_107) compile/launch failures in the CuTeDSL kernels #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 (Fix to #652: SM100 Grouped GEMM: read the TMEM accumulator before releasing the overlapping stage #654) and in SM107 grouped GEMM quant (Fix SM107 (Rubin) grouped GEMM quant: read TMEM accumulator before releasing the overlapping stage #673).
• Fixed the dGLU xfilter math (Fix bug in dglu xfilter math #568) and a missing attribute in Rubin grouped dGLU (fix missing attribute in Rubin grouped dGLU #625).
• Fixed missing engine class imports in the linear-attention tests (fix: GdnCuTileEngine import path in cutile conftest and _pygraph error msg #513) and a CI breakage from a missing is_exclusive guard (Fix ci breakage for missing protection to is_exclusive #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.

[coderabbitai]

Important

Review skipped

Auto reviews are disabled on base/target branches other than the default branch.

Please check the settings in the CodeRabbit UI or the .coderabbit.yaml file in this repository. To trigger a single review, invoke the @coderabbitai review command.

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 7dd80283-84f9-41cf-a025-071c0f9b9666

You can disable this status message by setting the reviews.review_status to false in the CodeRabbit configuration file.

Use the checkbox below for a quick retry:

• 🔍 Trigger review

---

_{Comment @coderabbitai help to get the list of available commands.}