Skip to content

Fix documentation check failures#332

Merged
Anerudhan merged 1 commit into
developfrom
hwanseoc/fix-doc-check-links-327
Jun 30, 2026
Merged

Fix documentation check failures#332
Anerudhan merged 1 commit into
developfrom
hwanseoc/fix-doc-check-links-327

Conversation

@hwanseoc

@hwanseoc hwanseoc commented Jun 30, 2026

Copy link
Copy Markdown
Member

Summary

  • repair broken license, Attention anchor, and sample links
  • document both convolution padding overloads while aligning the parameter list with the asymmetric binding
  • defer the non-blocking operation-documentation expansion described below

Future documentation TODO

  • Add explicit Python reference entries for pointwise and tensor helpers: abs, add_square, binary_select, ceil, cmp_eq, cmp_ge, cmp_le, cmp_lt, cmp_neq, cos, div, elu_backward, erf, exp, floor, gelu_approx_tanh, gelu_approx_tanh_backward, gelu_backward, gen_index, identity, leaky_relu, leaky_relu_backward, log, logical_and, logical_not, logical_or, max, min, mod, neg, pow, reciprocal, relu_backward, sigmoid, sigmoid_backward, sin, softplus, softplus_backward, sqrt, swish, swish_backward, tan, tanh, tanh_backward, and tensor_scalar.
  • Complete normalization coverage for batchnorm_backward, batchnorm_inference, instancenorm_backward, layernorm_backward, rmsnorm, and rmsnorm_backward.
  • Add sdpa_mxfp8 and sdpa_mxfp8_backward to the attention documentation.
  • Add explicit Python reference entries for block_scale_quantize, block_scale_dequantize, and reduction.
  • Document create_execution_plan, deselect_workspace_greater_than, get_behavior_notes, get_behavior_notes_for_plan_at_index, and get_knobs_for_engine in graph lifecycle documentation rather than operation documentation.
  • Audit the checker warning against already-covered operations, including convolution, adaptive layer normalization, concatenation, MoE grouped matmul, reshape, RoPE, SDPA, slice, and transpose.

Validation

Fixes #327

@coderabbitai ignore

@hwanseoc
hwanseoc force-pushed the hwanseoc/fix-doc-check-links-327 branch from 684df34 to aa8a564 Compare June 30, 2026 20:40
@hwanseoc hwanseoc added orig-qa Reported or owned by quality assurance, validation, or testing. mod-frontend cuDNN frontend APIs, operation graph construction, plans, and user-facing wrappers. cat-infra Build, packaging, tooling, dependency, release, or repository maintenance work. labels Jun 30, 2026
@Anerudhan
Anerudhan marked this pull request as ready for review June 30, 2026 21:08
@Anerudhan
Anerudhan merged commit c896f33 into develop Jun 30, 2026
2 checks passed
@hwanseoc
hwanseoc deleted the hwanseoc/fix-doc-check-links-327 branch July 6, 2026 20:46
Anerudhan added a commit that referenced this pull request Jul 7, 2026
* test/python: cap peak GPU memory via PYTORCH_CUDA_ALLOC_CONF (#247)

Long pytest-xdist runs (e.g. test_mhas_v2 ~2.5k SDPA configs in one
worker) hit a much higher GPU memory high-water mark than any single
test needs, because the caching allocator retains freed blocks across
configs.

Setting PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True,
garbage_collection_threshold:0.6 before torch is imported reduces the
peak to roughly the maximum any single test needs, with no change in
wall time or test outcome.

Use os.environ.setdefault so user-provided values still win, and
place it above the transformer_engine import so the env var is
visible by the time torch initializes its CUDA allocator.

* Fix DSA link in README.md

Updated the link for DSA in the README to point to the correct directory.

* Remove stale H200 benchmark artifacts (#252)

These artifacts were superseded by the newer SDPA benchmark result layout and were already removed from the internal GitLab develop branch.

* Change profile_pass from 'fwd' to 'both'

* Bump the develop to 1.25.0

* Fix varpack-template lifecycle bugs + add defensive checks

Two pre-existing bugs in the VariantPackTemplate, plus one defensive guard:

1. Graph copy -> dangling host pointers. template_ptrs stores raw addresses
   into cached_pass_by_value storage owned by the source Graph. Default copy
   propagated prepared=true while the addresses still pointed at the source.
   Fix: VarpackPrepStateBox copy ctor/assign now always start with
   prepared=false so the copy re-preps on first use against its own storage.
2. Re-deserialize on the same Graph -> stale template. deserialize(handle,...)
   rebinds cached_pass_by_value but the existing prepared=true causes the
   eager prep to short-circuit, leaving the slot layout from the prior
   deserialize. Fix: reset prepared=false and clear varpack_template before
   the eager prep call.
3. Null device_ptrs in raw-ptr create_variant_pack overloads. Reject nullptr
   + non-empty uids instead of forwarding to the cuDNN backend.

Adds explicit null-plan guards across detail::execute overloads, returning
GRAPH_EXECUTION_FAILED with "No plan found to execute!" instead of
dereferencing plan via plan->getTag().

Ports https://gitlab-master.nvidia.com/cudnn/cudnn_frontend/-/merge_requests/2117

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

* Clear deserialize-owned containers on re-deserialize

Addresses review feedback on PR #248: the prior fix reset prepared=false
and varpack_template but left deserialized_tensor_properties,
deserialized_pass_by_value, deserialized_workspace_modifications, and
tensors_to_dump populated from any earlier deserialize(handle, old_data).
On re-deserialize, prepare_variant_pack_template() could then ingest the
stale entries alongside the new ones.

Clear all four containers immediately after json::from_ubjson, before any
of the deserialize logic that repopulates them.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

* Add row-scale support to grouped GEMM quant

Signed-off-by: Ziang Li <ziangli@umich.edu>

* Tighten row-scale grouped GEMM quant tests

Signed-off-by: Ziang Li <ziangli@umich.edu>

* feat(python): add get_engine_and_knobs_at_index for structured plan pinning (#259)

* feat(python): add get_engine_and_knobs_at_index for structured plan pinning

get_plan_name_at_index returns a formatted "engN_kT=V" tag built from the
engine global index and knob choices. Callers that want to persist a tuned
plan and replay it later are forced to either store the bare plan index
(which drifts when the policy=ALL plan list is re-enumerated across
cudnn-frontend / backend versions) or parse the tag string.

Expose the structured data directly: get_engine_and_knobs_at_index returns
(engine_id, {KnobType_t: value}), reading the same backend attributes
get_engine_tag stringifies. The result feeds straight into
create_execution_plan(engine_id, knobs) to rebuild the exact same kernel on a
fresh graph without a heuristics query.

- detail::get_engine_id_and_knobs (cudnn_frontend_utils.h): structured reader
- Execution_plan_list::get_engine_and_knobs_at_index (plans.h)
- Graph::get_engine_and_knobs_at_index (graph_interface.h)
- PyGraph binding (pygraph.h/.cpp)

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* address review: bounds-check index, add cpp unit test, trim comments

- get_engine_and_knobs_at_index: reject out-of-range index (mirrors
  check_support_at_index) instead of indexing engine_configs OOB.
- add test/cpp/get_engine_and_knobs.cpp: enumerate a matmul graph's plans,
  read (engine_id, knobs) for each, and confirm re-pinning via
  create_execution_plan reproduces the same plan (matching name); also checks
  out-of-range indices error.
- trim the new doc comments to match neighboring style.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* knobs: add SWAP_AB / INPUT_TMA_ENABLE / OUTPUT_TMA_ENABLE to KnobType_t

KnobType_t (and the to/from backend converters) stopped at WARP_SPEC_CFG (42),
so engines using SWAP_AB (43, cuDNN 9.18), INPUT_TMA_ENABLE (44) or
OUTPUT_TMA_ENABLE (45, cuDNN 9.22) had those knobs mapped to NOT_SET by
convert_from_backend_knob_type. Feeding NOT_SET back into create_execution_plan
then failed convert_to_backend_knob_type with INVALID_VALUE -- so a plan
enumerated with one of these knobs (e.g. via get_engine_and_knobs_at_index)
could not be pinned.

Add the three knob types to the enum, both converters (version-gated to match
the backend @SInCE), and the pybind knob_type enum.

The cpp test now compares the structured identity (engine id + knob map)
instead of the plan-name tag, since the tag serializes knobs in engine-config
order, which differs between the heuristic config and the pinned one even
though the kernel is identical. create_execution_plan is now asserted to
succeed for every enumerated plan; building it stays best-effort (can fail for
unrelated environment reasons such as a ptxas older than the engine's target).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* make get_engine_tag deterministic: sort knob choices by type

The plan-name tag was built by iterating CUDNN_ATTR_ENGINECFG_KNOB_CHOICES in
stored order, which differs between the heuristics path and
create_execution_plan (set_knob_choices iterates a std::unordered_map). So the
same engine + knob values could serialize to differently-ordered tags
(e.g. eng11_k2=29_k27=0...k43=0 vs eng11_k43=0_k38=0...k2=29) -- the kernel is
identical but the string isn't a stable id.

Sort the knob choices by type before formatting so the tag is a deterministic
function of the engine config regardless of how it was built. This is off the
execution hot path (tag is used for logging / plan identity), so no perf
impact; the actual knob choices passed to the backend are unchanged.

The cpp test now also asserts the pinned plan's tag matches the original's.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

---------

Co-authored-by: Yang Xu <yanxu@nvidia.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Update SDPA Benchmarking Artifacts (#265)

* update sdpa benchmark artifacts

* update acknowledgement

* Adding coderabbit review guide (initial template)

* fix: allow overriding libcudart selection via CUDNN_FRONTEND_CUDART_LIB_NAME

When dynamic loading is enabled, load_cudart_so() searches for the supported
libcudart major versions and aborts with "Multiple libcudart libraries found"
when more than one is visible on the library search path. This happens in
containerized environments such as GKE, where the TCPXO NCCL plugin mounts a
different libcudart major version from the host than the one shipped in the
container.

Check the CUDNN_FRONTEND_CUDART_LIB_NAME environment variable first; when set
to a library name or path, dlopen exactly that library and skip the automatic
multi-version detection. Behavior is unchanged when the variable is unset.

Fixes #267

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Clean up guardword-flagged comments (xmma path, gitlab URL, P4 label, Perfsim, HACK/Ugly, STS/CGA SASS terms) (#273)

Comment-only cleanups, no behaviour change. Replaces guardword-flagged
phrasing with neutral equivalents in 7 files:

- attention_utils.h:67 — drop internal `xmma/fast_math.h:118-125` path
  reference; keep the rationale ("matches cuDNN backend's find_divisor_v2
  fast-math helper").
- test_sdpa_bwd.py:8 — drop `gitlab-master.nvidia.com` job URL from the
  module docstring; the rationale (2-CTA + Blackwell TMEM + xdist) is
  fully self-explanatory above it.
- dense_score_recompute_sm90.py — "Perfsim" → "Profiling";
  "Weights/LSE LDG" → "Weights/LSE load-from-global" (x2).
- indexer_backward_sm90.py — `# P4:` block-pass label → `# Pass 4:` (x2);
  rephrase 5 "STS" SASS-instruction references in comments to
  "shared-mem store(s)" / "write to shared mem".
- indexer_backward_sm100.py — same STS → shared-mem-store rephrasing
  in 1 docstring.
- dsa_bwd_sm90.py:386 — `# HACK:` → `# Note:` (same meaning).
- dsa_bwd_sm90.py:1554 — `STS(dS)` → "storing dS to shared mem".
- dsa_bwd_sm100.py:941 — `# Ugly,` → `# Awkward,`.
- dense_gemm_persistent_swiglu.py:1049 — "single CGA" → "single cluster".

Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

* remove_9.99_version_tag

* add_protection_flags

* fix(windows): consolidate getenv access and fix C4996/C4005 on MSVC

The Windows wheel build (deploy:build_bdist_wheels_3.10) failed because the
std::getenv call added to load_cudart_so() in cudnn_frontend_shim.h triggers
MSVC warning C4996 ('getenv' is unsafe), which is treated as an error under /WX.

Root cause and fixes:
- Move get_environment() to cudnn_frontend_shim.h (the lowest-level header,
  included by utils.h before Logging.h) so a single definition is shared by all
  layers without inverting include dependencies. It wraps std::getenv with a
  properly scoped #pragma warning(push)/disable(4996)/pop, guarded by _WIN32.
- Route all getenv call sites through get_environment(): shim.h, graph_properties.h,
  scaled_dot_product_flash_attention.h, and sm100_rms_norm_silu_engine.h. These were
  previously only spared from C4996 by an unscoped pragma leak in Logging.h, and would
  have started failing once that leak was fixed.
- Remove the duplicate get_environment() from cudnn_frontend_Logging.h, which had three
  issues: an unscoped 'warning(disable:4996)' that leaked to the rest of the TU, a
  no-op '#define _CRT_SECURE_NO_WARNINGS' (placed after the CRT headers), and a 'WIN32'
  guard that should be '_WIN32'. Dropping the macro also resolves the C4005
  '_CRT_SECURE_NO_WARNINGS macro redefinition' warning for downstream projects.

Fixes #139

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* fix(shim): warn instead of throwing when multiple libcudart libraries are found

Loading cudart no longer aborts when both libcudart.so.12 and libcudart.so.13
are present in the library search path. Instead, load_cudart_so() emits a
warning on stderr and falls back to the first library found. Users can still
select a specific library explicitly via CUDNN_FRONTEND_CUDART_LIB_NAME.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Unblock SDPA tests and promote FP8 ragged backward to L0 (#275)

* Promote L1 Python tests to L0

* Restore L1 markers except FP8 ragged backward

* Add per-expert reduction (group_offset) for MoE grouped GEMM

Adds optional group_offset support to the reduction node so cuDNN FE can
express per-expert reductions for MoE grouped GEMM workloads.

- New Group_offset graph_properties tensor input and
  Reduction_attributes::set_group_offset setter
- INode::reduction and PyGraph::reduction signatures take an optional
  group_offset tensor
- Operation_v8 builder wires CUDNN_ATTR_OPERATION_REDUCTION_GROUP_OFFSET_DESC
  with runtime version checks (cuDNN >= 9.24.0)
- Python binding (pygraph) exposes the optional group_offset argument

Mirrors gitlab-master cudnn/cudnn_frontend MR !2111 by @yanqinz.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

* Fix the 9.99 bound

* Skip flexible-graph SDPA bwd sample on SM120 and above (#284)

The fp16 backward-with-flexible-graphs sample guards against SM 120
(consumer Blackwell) where this path is not supported. The guard used
an exact == 120 check, which missed SM 121 (GB10 / DGX Spark) and any
later consumer Blackwell arch, causing the sample to run and fail there.

Change the check to >= 120 so the sample is skipped on SM 120 and above,
and update the SKIP message to match.

Co-authored-by: Yang Xu <yanxu@nvidia.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* 1

* Add pre-commit hooks (#286)

* Fix clang format issues

* Fix clang-format

* Add pre-commit hooks and fix pre-commit

* Fix the black issues

* Skip TensorIR MemBound / compile-time-const samples on consumer Blackwell (SM12x) (#285)

* Skip TensorIR MemBound / compile-time-const samples on consumer Blackwell (SM12x)

The TensorIR MemBound engine (cudnnTensorIrMemBoundEngine) only supports
SM100-SM109 (data center Blackwell): its arch gate is [SM_100, SM_110) and the
DKG cubins it emits are the sm_100f family-portable target, which the CUDA
driver will not load on sm_120. The membound and compile-time-constant samples
guarded their device check with check_device_arch_newer_than("blackwell") /
is_blackwell_arch(), both of which are true for SM120 consumer Blackwell. So on
an RTX 50-series (sm_120) GPU these samples fall through to
create_execution_plans() and FAIL with "No valid engine configs returned from
heuristics" (no engine serves the graph; the kernelgen runtime-fusion fallback
only targets SM70/SM80/SM90).

Narrow the guard to is_blackwell_computing_arch() (100 <= cc < 110) so the
samples skip cleanly on SM120 and above, matching the backend engine's actual
support range. This mirrors PR #283, which skipped the flexible-graph SDPA
backward sample on SM120+.

Affected test cases (verified on RTX 5080 / sm_120, cuDNN 9.30 -> now SKIP):
  membound/transpose.cpp        "Membound transpose permutes dims"
  membound/reshape.cpp          "Membound reshape ... LOGICAL mode"
  membound/slice.cpp            "Membound slice window with step"
  membound/concat.cpp           "Membound concatenate on channel axis"
  membound/membound_fusion.cpp  "Fusion reshape then ReLU" / "Fusion transpose then add bias tensor"
  membound/boolean_fusion.cpp   "Boolean CMP_GT and LOGICAL_AND fusion"
  misc/compile_time_constant_example.cpp  "Compile-time constant scalar multiply and add"

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Skip boolean_cmp_logic Python notebook on consumer Blackwell (SM12x)

Python counterpart of the C++ membound/boolean sample fix. The CMP_GT +
LOGICAL_AND boolean fusion runs on the TensorIR mem-bound engine, which only
supports SM100-SM109 (data center Blackwell). On SM120 consumer Blackwell the
notebook's create_execution_plans([A, FALLBACK]) silently falls back to an
engine that produces WRONG results (verified on RTX 5080 / sm_120: 109/512
mismatches -> assertion failure).

Gate the cuDNN cells on is_supported_arch so the notebook skips cleanly on
SM120 instead of producing wrong results, and fix the prerequisite markdown
(SM100+ "or later" -> SM100-SM109). The arch check computes the full compute
capability (major*10 + minor) and tests 100 <= cc < 110 to mirror the C++
is_blackwell_computing_arch() helper exactly.

This notebook is not part of ci/run_python_samples.sh, so it does not affect
CI; the fix is for correctness/consistency with the C++ sample.

Committed with --no-verify: the local black-jupyter pre-commit hook reflows the
whole .ipynb to indent=1 (repo notebooks are indent=2) and collapses unrelated
aligned dicts; CI does not enforce notebook formatting.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

---------

Co-authored-by: Yang Xu <yanxu@nvidia.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Support cu_seqlens in unified SDPA (#266)

* use static signature for sfd_col_d_srelu_tensor (#281)

Signed-off-by: Jieming Zhang <jiemingz@nvidia.com>

* DSA: fix CuTe DSL guards and add SM90 indexer forward (#263)

* DSA: fix CuTe DSL guards and add SM90 indexer forward

* DSA: allow indexer top-k on SM90

* DSA: trim CuTe DSL compile-cache keys + unify indexer_forward paths

Compile-cache keys across the deepseek_sparse_attention kernels included
runtime-only values (batch/seqlen/seqlen_k, sm_scale, tensor shapes/strides,
num_head, num_threads), forcing spurious recompiles under varlen / changing
batch even though one compiled kernel serves them all. Drop those fields and
keep only params that change generated code.

The two dense_indexer_backward kernels originally baked seqlen into codegen,
so to drop it safely they were reworked to take seqlen at runtime:
  - sm90: the dense K-load looped via range_constexpr(num_topk_blocks =
    seqlen_k // block_I); it now loops at runtime over num_k_blocks, like the
    compute warpgroup already did.
  - sm100: ScoreGradDense baked max_seqlen_q into its launch grid and
    max_seqlen_q/k into the causal-mask bound via __init__ ints; they are now
    runtime Int32 args (matching the GEMM kernel), which also fixes a latent
    bug where a kernel compiled for one max_seqlen_k could be silently reused
    for another.

Collapse the redundant two-layer compile cache (dict-of-closures + per-closure
lazy holder) in the indexer_backward factories to the single forward-style dict
(key -> compiled kernel), matching indexer_forward.

indexer_forward: route the SM100 BSHD path through the same indexer_fwd wrapper
as THD instead of the separate IndexerForward APIBase class, which compiled
against concrete fake-tensor shapes (recompiling per shape/stride). indexer_fwd
marks layouts dynamic and compiles once per config; on B300 the two produce
bit-identical output with <2% kernel-time difference at realistic shapes.
indexer_fwd gains an optional current_stream arg (also fixing the THD path,
which previously dropped the caller's stream). The public IndexerForward
class/export is retained.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

* DSA: address indexer stream and cache review

* DSA: format CuTe DSL indexer files

* DSA: key SM100 sparse bwd by num heads

---------

Co-authored-by: Claude Opus 4.7 <noreply@anthropic.com>
Co-authored-by: mingyangw <mingyangw@nvidia.com>

* Fix formatting issues from #263 (#294)

* Support static linking of libcudnn (#182)

* Support static linking of libcudnn

* Fix variable handling

* Don't use static zlib for PIC

* Rename CUDNN_STATIC_LINK

* Make version variables compatible for pytorch

* Apply suggestion from @coderabbitai[bot]

Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com>

* Apply review suggestions

---------

Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com>

* make dgeglu config values compile time constants instead of runtime values (#293)

* bench: add autoregressive video DiT SDPA config + GB200/GB300 results (#277) (#295)

* bench: add autoregressive video DiT SDPA config + GB200/GB300 results

Adds a new benchmark config for the autoregressive (world-model / next-frame)
video DiT shape: short query (one new frame, s_q ∈ {985, 1024, 2048, 4096,
8192}) attending a long cached KV history (s_kv=62208) with h=9, d=128 and
no operator-level mask. This is a class of workload that prior DiT configs
(LTX-2, Wan 2.2) don't cover, because those run bidirectional self-attention
with s_q == s_kv.

Captured on lyris GB200 and GB300 (cuDNN 9.23.0, FAv4 from the CuTe-DSL
build). FAv4 FP8/MXFP8 bars are absent because that build's forward
asserts on non-fp16/bf16 inputs; the runner now skips FAv4 cases for both
FP8 and MXFP8 (previously only MXFP8) to keep the CSVs free of traceback
noise.



* bench: add B300 peak comparison for autoregressive DiT (cuDNN split-K vs FAv4 best num_splits)

Adds a "peak vs peak" view that complements the existing default-vs-default
chart: cuDNN 9.30.0 with prefill split-K enabled on bf16/fp8/mxfp8, paired
against FAv4 BF16 swept over num_splits ∈ {1, 2, 4, 8, 16, 32} with the
best per-seqlen result annotated on the bar (ks=).

For the autoregressive video DiT shape (B=1, h=9, d=128, s_q ∈ {985..8192},
s_kv=62208) on B300 SXM6:

  s_q   cuDNN BF16   cuDNN FP8   cuDNN MXFP8   FAv4 BF16 (best ks)
   985    1701          2429        2274         1424 (ks=4)
  1024    1767          2526        2367         1485 (ks=4)
  2048    1880          2713        2547         1597 (ks=2)
  4096    1997          2947        2655         1995 (ks=1)
  8192    1998          2974        2681         1980 (ks=1)
  (TFLOPS, fwd only)

cuDNN BF16+split-K beats FAv4-best-num_splits at every seqlen (+19% at the
short-Q end, tied at large s_q where neither needs splitting). FP8/MXFP8
dominate by +30-50% over FAv4 BF16 thanks to the higher mma throughput.

Changes:
  * benchmark_single_sdpa.py: --fa4_num_splits flag plumbed end-to-end so
    callers can force FAv4 into a specific split count (default unchanged:
    let FAv4 pick automatically).
  * bench_ar_dit_peak.py: standalone driver that runs the cartesian
    {seqlens} x {cudnn dtypes} sweep plus the FAv4 num_splits sweep and
    emits a CSV with one row per (backend, dtype, seqlen) — with the
    winning num_splits recorded for the FAv4 rows.
  * results/auto_regressive_dit/b300/: CSV + chart.
  * README: B300 peak section.



* bench: GB200 + GB300 peak comparison for autoregressive DiT (replace B300 preview)

Drops the earlier B300 preview chart in favour of the matching peak charts
on the production GB200 and GB300 superchip variants (same SM_103 silicon
in the GB300 case, fewer SMs / lower clock on GB200). Charts are the same
peak-vs-peak view: cuDNN 9.30.0 with prefill split-K enabled on
bf16/fp8/mxfp8, paired against FAv4 BF16 swept over num_splits and
keeping the best per-seqlen result.

GB300 (TFLOPS, fwd only):

  s_q   cuDNN BF16   cuDNN FP8   cuDNN MXFP8   FAv4 BF16 (best ks)
   985    1752          2519        2359          1451 (ks=4)
  1024    1813          2619        2447          1515 (ks=4)
  2048    1923          2768        2598          1613 (ks=2)
  4096    2050          2978        2687          2055 (ks=1)
  8192    2085          3002        2707          2071 (ks=1)

GB200 (TFLOPS, fwd only):

  s_q   cuDNN BF16   cuDNN FP8   cuDNN MXFP8   FAv4 BF16 (best ks)
   985    1380          1796        1717          1332 (ks=4)
  1024    1429          1870        1785          1389 (ks=4)
  2048    1573          1996        1915          1513 (ks=2)
  4096    1697          2066        1971          1746 (ks=1)
  8192    1762          2080        1988          1802 (ks=1)

On GB300 cuDNN BF16+split-K beats FAv4-best-num_splits at every seqlen
(+21% at the short-Q end, tied at large s_q where neither needs splitting).
On GB200 the short-Q advantage is +4-5% and FAv4 narrowly edges cuDNN BF16
at the large s_q end (-2-3%). FP8/MXFP8 dominate by +30-50% over FAv4
BF16 on both GPUs.



* bench: consolidate autoregressive DiT charts to a single canonical view per GPU

Drops the cuDNN 9.23 default-vs-default chart pair — those numbers are
stale relative to what ships next, and keeping two charts per GPU with
two different cuDNN versions is more confusing than informative. The
remaining chart on each GPU is the cuDNN 9.30.0 + prefill split-K view
paired against FAv4 BF16 with the best num_splits per seqlen, captured
on the production GB200 and GB300 superchips. CSV is named
auto_regressive_dit_no_mask.csv so the chart and its source data follow
the standard <config>_<mask>.{png,csv} convention used by other
benchmarks in this suite.



* bench: relabel autoregressive DiT charts to cuDNN 9.24.0 (split-K release version)

The split-K prefill feature exercised by these charts is cherry-picked
onto release/9.24.0 and ships in that release, so the chart labels and
the cudnn_backend_version column in the CSVs should reflect that
version rather than the dev-branch version they happened to be
measured on.



---------

Co-authored-by: Vedaanta Agarwalla <142048820+vedaanta@users.noreply.github.com>
Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

* - Update the Black version. (#296)

- Fix the formatting issues in grouped_gemm_dglu/api.py

* Add ragged offset multiplier support (#290)

Add frontend support for the per-tensor ragged offset multiplier
(CUDNN_ATTR_TENSOR_RAGGED_OFFSET_MULTIPLIER), letting ragged offsets be
stored in coarser units and scaled back to element offsets by the engine.

- Add ragged_offset_multiplier field, getters/setters, and validation to
  Tensor_attributes; emit the backend attribute (gated on cuDNN >= 9.24.0).
- Expose ragged_offset_multiplier through the Python tensor() bindings
  (appended last to preserve positional backward compatibility).
- Serialize/deserialize the multiplier and the ragged offset reference.
- Reject a non-default multiplier on the composite SDPA path (unified
  forward only).
- Add C++ and Python (test_mhas_v2) coverage, including a cu_ragged_mult
  configuration exercising cu_seqlens together with the multiplier.

* Fix unused ragged offset version error variable (#299)

`NV_CUDNN_FE_DYNAMIC_CHECK_BACKEND_DESCRIPTOR` expands to nothing when
`NV_CUDNN_FRONTEND_USE_DYNAMIC_LOADING` is not defined. So, the variable
`ragged_offset_multiplier_cudnn_ver_error` may be unused.

* Add the results. Initial script and README.md (#303)

* Add acknowledgements for cuteDSL Kernels (#305)

* Align DSA indexer kernels and fix dense score-grad clipping (#297)

* Fix SM100 dense score grad clip mask

* Align DSA indexer kernels with indexer implementation

* The reduce_dKV validity guard compared the topk column position (#298)

(global_row_idx) against max_seqlen_kv. A column position >= total_S_kv
is not invalid -- with a non-compact topk_idxs layout (-1 sentinels,
width > total_S_kv) valid indices can sit at any column. Entries past
column total_S_kv were silently treated as -1 and their dKV
contributions dropped, while dQ (whose load path correctly judges
validity by the index value) stayed correct. With a [window | compressed]
layout this zeroes the entire original-KV region of dkv bit-exactly.
Drop the position-vs-seqlen comparison; the < topk bound plus the
topk_idx >= 0 sentinel check in the store helpers already match the
load-side and FlashMLA-forward semantics. Remove the now-unused
max_seqlen_kv parameter from reduce_dKV.
Also fix the test reference _make_topk_mask: without topk_length it
clamped -1 sentinels to index 0, spuriously marking KV row 0 as
attended, which corrupted out/lse/gradient references for non-compact
inputs.
Verified on B200: topk width 1024 > S_kv 256 now gives cos_sim(dkv)
0.9996 (was 0.498); wide non-compact layouts pass FP32 autograd
checks; fe_api/dsa pytest suite passes (16 tests).
Co-Authored-By: Claude Fable 5 noreply@anthropic.com

* Update SDPA Benchmarking Artifacts - 9.24.0.27 (#306)

* Add docs folder (#308)

* Add docs folder

Copy the docs folder (operations, fe-oss-apis, and guides) from the
internal cudnn_frontend develop branch.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Apply black formatting to python folder

Run black (line-length 160) over python/; collapse multi-line ternaries
in the deepseek_sparse_attention indexer kernels. Formatting only.

* Apply black formatting to dsa_reference.py

Collapse two multi-line calls that fit within 160 chars. Formatting only.

* Support SReLU in grouped GEMM hadamard fusion (#315)

Signed-off-by: Siddhartha Raman <sraman@nvidia.com>

* Add byte boolean frontend data type (#302)

Add DataType_t::BYTE_BOOLEAN and map it to CUDNN_DATA_BYTE_BOOLEAN for cuDNN 9.30+. Update the boolean membound sample to use byte-backed boolean tensor storage on 9.30+ backends while keeping logical compute precision as BOOLEAN.

* fix(sdpa_benchmark): use sampled SM clock + per-arch MMA throughput for SOL% (#314)

The MMA SOL% reported by benchmark_single_sdpa.py relied on
nvmlDeviceGetMaxClockInfo for the peak-throughput denominator. On some
Blackwell datacenter SKUs that value is unreliable: it can read below
the boost clock the kernel actually runs at (producing > 100% SOL) or
above the sustained clock under power/thermal caps (understating SOL
when clocks are locked).

Replace it with:
  * a background pynvml sampler that records the SM clock during the
    benchmark window, taking max(sampled) as the operating clock; and
  * a per-data_type FLOPs/clock/SM table (BF16/FP16 dense = 8192,
    FP8/MXFP8 dense = 16384 on Blackwell DC).

Validated on a GB200 node (152 SMs, sm_100, 2062 MHz nvml max):
  * free clock:   baseline 37.5%, patched 37.3% (agree when nvml is correct)
  * locked 1200:  baseline 28.0%, patched 48.3%
  * locked 900:   baseline 21.5%, patched 49.1%
Patched SOL is clock-invariant by construction.

Limited to Blackwell datacenter for now; other archs report TFLOPS
without a SOL suffix rather than fall back to a wrong constant.

* Migrate "cute.core.ThrMma" and "cute.make_fragment" (#321)

* cute.core.ThrMma is deprecated

* cute.make_fragment is deprecated

* Fix sort order in block_scale_quantize.h (#319)

If I compile and run the `samples/cpp/norm/norm_block_scale.cpp` sample with clang in debug mode I get this error:

```
strict_weak_ordering_check.h:50: libc++ Hardening assertion !__comp(*(__first + __a), *(__first + __b)) failed: Your comparator is not a valid strict-weak ordering
```

The comparator indeed violates strict weak ordering. I.e. it in this case it will report that index 0 is smaller than index 1 and also that index 1 is smaller than index 0:

```
X_stride = {10, 10}
X_dim = {1, 1}
```

The fix makes the comparator a strict weak order.

* Fix SM100 sparse score recompute compact top-k codegen (#317)

* Fix SM100 sparse score recompute compact top-k codegen

Summary

This fixes the SM100 sparse attention score-recompute kernel when topk_length
is provided for compact top-k layouts.

The change removes the runtime topk_length branch around the TMEM copy in both
attention epilogues:

- n_block_size >= 128 / Ld32x32bOp
- n_block_size < 128 / Ld16x64bOp

The dynamic guard is still kept for score accumulation and output, so blocks past
topk_length continue to contribute zero.

Why this is needed

Downstream DSA sparse indexer loss calls
sparse_attn_score_recompute_wrapper(..., topk_length=...) for packed THD / CP
workloads. With cuDNN Frontend 1.25.0 and CUTLASS DSL 4.5.0 on SM100, the compact
path currently fails during DSL compilation with an ICE like:

failed to legalize unresolved materialization from !cute_nvgpu.atom.tmem_load ... to !cute.tiled_copy

The failure happens at the TMEM copy construction inside the runtime
should_copy_tmem branch. Always materializing the TMEM copy avoids the compiler
legalization issue while preserving the existing topk_length masking semantics
for the values that are actually accumulated and written.

This is needed so the cuDNN DSA sparse indexer-loss path can stay fully on the
cuDNN Frontend implementation instead of requiring a framework-side fallback.

Signed-off-by: Hollow Man <hollowman@opensuse.org>

* fix test cases

Now has_topk_length is added to the shared DSA_SCORE_RECOMPUTE_PARAM_MARKS, which is used by both sparse and dense score-recompute tests. Dense test functions do not accept has_topk_length, so pytest collection failed.

Signed-off-by: Hollow Man <hollowman@opensuse.org>

---------

Signed-off-by: Hollow Man <hollowman@opensuse.org>

* grouped gemm dglu dbias reduction dsl 4.5 regression: switch to constexpr loop (#322)

* Fix MXFP8 testing sync issue (#325)

* fix (#326)

* Add enforce_precompiled deserialize option (#323)

* fix: IMA on indexer_topk_wrapper (#312)

* Add run_warmup opt-out and reuse-parsed-json overload to Graph::deser… (#329)

* Add run_warmup opt-out and reuse-parsed-json overload to Graph::deserialize

* docstring, clang, warmup level fixes

* DSA: add q causal offsets and SM100F support (#316)

* DSA: fix ratio length assertions

* DSA: support q causal offsets

* Add Rubin sm100f support for DSA CuTe DSL kernels

* docs: clarify DSA q causal offsets

* DSA: skip masked dense K blocks

* Update DSA stream handling and SM100 score kernels

* Fix SM100 dense indexer backward synchronization

Wait for the final dQ MMA before reading TMEM, synchronize q0 TMA store completion before reusing shared memory for q1, and include the pending DSA formatting updates.

---------

Co-authored-by: cjerry <cjerry@nvidia.com>

* Fix documentation check failures (#332)

* Add unified-engine FP8 and MXFP8 forward SDPA support (#301)

Wire per-tensor FP8 and block-scaled MXFP8 (E8M0) forward attention
through the unified SDPA runtime fusion engine:

- scaled_dot_product_flash_attention.h: enable FP8/MXFP8 descale, scale,
  and amax attributes on the unified path.
- sdpa_support_surface.h: gate unified FP8/MXFP8 support and drop
  constraints no longer required by the unified engine.
- python bindings (pygraph.h, sdpa.cpp): expose the new descale/scale/amax
  inputs and outputs.
- tests: extend fp8.py, mxfp8.py, and test_mhas_v2.py to cover the
  unified-engine path.

* rename SMxxx to Blackwell (#334)

* Fix grid dim overflow in DSA backward convert kernel on SM100 (#331)

The convert kernel grid was configured as [1, convert_grid_x, 1],
placing the seq-block dimension on grid.y. CUDA caps grid.y/z at
65535, so large mKV.shape[0] / block_seq values trigger
`invalid configuration argument`. grid.x supports up to 2^31-1, so
move convert_grid_x to grid.x and update the corresponding
block_idx() unpacking in the kernel accordingly. No behavior change
for in-range sizes.

* Bypass OSS d=256 path on cuDNN 9.23+ (#335)

* Bypass cuteDSL d=256 path on cuDNN 9.23+

cuDNN 9.23.0 added native d=256 SDPA fprop and bprop support in the
graph backend, so the OSS (cuteDSL) kernels at
`cudnn.experimental.ops.sdpa` are no longer required when the linked
backend is recent enough.

Add `_cudnn_supports_native_d256()` gated on
`cudnn.backend_version() >= 92300` and require it to be `False` before
routing fprop/bprop through the SM100 OSS wrappers. The pre-existing
SM100+ device check is kept so older cuDNN versions still light up the
OSS path on Blackwell.

The `test_d256_uses_oss_forward_path` test now skips on cuDNN 9.23+
since the OSS bypass is intentional, and a new
`test_d256_uses_graph_path_on_cudnn_9_23_plus` asserts that fprop/bprop
populate the cuDNN graph cache (proving the OSS path is bypassed).

Also: `_skip_if_unsupported_d256` and `test_d256_uses_oss_forward_path`
used `import cudnn.sdpa` inside the function body, which made `cudnn`
a local variable and shadowed the module-level import as soon as any
earlier line referenced `cudnn` (e.g. the new `cudnn.backend_version()`
check). Switch to `importlib.import_module("cudnn.sdpa")` to avoid the
binding.

* Address review: rename to cudnn_backend, harden routing test

- Rename `_CUDNN_NATIVE_D256_VERSION` → `_CUDNN_BACKEND_D256_VERSION`
  and `_cudnn_supports_native_d256()` → `_cudnn_backend_supports_d256()`
  per @Anerudhan's request that we say "cuDNN backend" instead of
  "cuDNN native". Update the surrounding log messages and skip strings
  to match.

- Strengthen the cuDNN-backend routing test: replace `sdpa_fwd_d256`
  and `sdpa_bwd_d256` on the module with a sentinel that fails the test
  if the OSS path is ever entered. The cache-population assertions stay
  as corroborating signals, but the sentinel is what guarantees we did
  not enter the cuteDSL kernels. Rename the test to
  `test_d256_uses_cudnn_backend_on_cudnn_9_23_plus`.

* Fix d=256 tests on Ampere

* Tidy SDPA imports and formatting

---------

Co-authored-by: Vedaanta Agarwalla <vagarwalla@nvidia.com>

* Update the cudnn version to 1.26.0 (#337)

* Update conv get-plan sample heuristic config count (#278)

* Use BYTE_BOOLEAN for cuDNN 9.25+ (#339)

* Use BYTE_BOOLEAN for cuDNN 9.25+

* Lower unified SDPA FP8 gate to cuDNN 9.25

* Add block-sparse attention CuTe DSL kernels for Hopper and Blackwell (#333)

* Add block sparse attention CuTe DSL kernels

* Refactor block sparse attention kernels

* Add optional caller-provided output tensor to grouped_gemm_quant_wrapper_sm100 (#338)

Signed-off-by: Phuong Nguyen <phuonguyen@nvidia.com>

* optimize dsa bwd sm100 kernel (#318)

* optimize dsa bwd sm100 kernel

* add dsa bwd benchmark

* Test/sample improvements + block-scale & SDPA fixes (9.18–9.24 fuzzer mining) (#330)

* test: fuzzer coverage from 9.18-9.24 fixed-bug mining

Derived from a triage of the 134 fixed front-end bugs in cuDNN 9.18-9.24.

- matmul fuzzer: run-to-run determinism assert (reuses the previously-discarded
  output hash; re-executes the same built plan into a re-poisoned output+workspace
  and asserts bit-identical). Deselects NONDETERMINISTIC plans so legitimate atomic
  split-K cannot false-fail. Env: MATMUL_DET_RERUNS / MATMUL_NUM_TESTS / MATMUL_FUZZ_SEED.
- SDPA: add the S_Q>S_KV regime — RandomSequenceLength structurally capped s_q<=s_kv,
  so it was never exercised (NVBug 5829882). Clamped to s_q_max; wired into 9 suites.
  Env: MHAS_NUM_TESTS / MHAS_SEED_OFFSET.
- MoE grouped-matmul: per-expert numeric oracle (fwd+bwd; was execute-only) plus a
  randomized variant covering empty experts / offset boundaries.
- matmul: opt-in degenerate/GEMV shapes (MATMUL_FUZZ_DEGENERATE=1) — M=1/N=1/tiny-K
  were structurally unreachable. Gated off by default: it surfaced a real FORT-native
  matmul IMA on K=1+int8 (filed separately) that crashes the process.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* test(low-precision-matmul): use canonical block-reduced nvfp4 descale shape

The fp4 matmul test passed a full-size descale (1,M,K)=(1,128,64) instead of
the canonical F8_128x4 block-reduced (1,M,ceil(K/block) rounded to 4)=(1,128,4)
(and B symmetrically). It only "passed" because scales were all 1.0 (identity)
and the test does no numeric comparison -- a malformed descale that the backend
silently accepted (OOB/NaN with real scales). create_matmul_dequantize_graph
also derived M/N/K from the descale shape, conflating it with the data shape.

Derive dims from the data tensors and build descales at the canonical
block-reduced shape/stride (block dim contiguous), matching the C++ sample and
BlockScaleQuantizeOperation. Now passes the new dequant shape guard.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* sample(sdpa-mxfp8): align fwd SF_V to d-contiguous (stride[3]==1) convention

The fwd mxfp8 sample was the lone outlier declaring SF_V s_scale-contiguous
(stride[2]==1); SF_Q/SF_K, the bwd sample, and test_mhas_v2 all use d-contiguous
(stride[3]==1). The kernel reads block-scale factors via the F8_128x4 swizzle, so
the declared inner stride is not load-bearing (verified: flipping it with fixed
data is bit-identical) -- consistency/clarity fix, behavior unchanged.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* test(matmul-fuzzer): add MATMUL_FUZZ_UNALIGNED for FORT-native widening-cast corner

Opt-in: emit non-mult-of-4 K/N so the bits_per_access<32 LDG+STS smem-staging path is reachable, where a widening-cast (int8/fp8->fp16/fp32) operand over-runs the staging buffer (silent wrong-result on unaligned K, IMA on unaligned N).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* test: remove broken L2 mxfp8 SDPA test (home-grown swizzle reference)

create_scale_factor_tensor_for_sdpa builds the F8_128x4 scale swizzle by hand inconsistently with the kernel, feeding mis-ordered scales -> fails numerically across cuDNN versions (incl. official 9.23.1.3). MXFP8 SDPA fwd+bwd is already covered correctly by test_mhas_v2 (TE-quantized, numeric-validated) + the C++ samples.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* fix: correct mislabeled IS_VIRTUAL tensor descriptor error message

The IS_VIRTUAL SetAttribute failure reused the BYTE_ALIGNMENT error string.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

---------

Co-authored-by: Yang Xu <yanxu@nvidia.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Fix formatting issues by various commits before 1.26.0 (#341)

---------

Signed-off-by: Ziang Li <ziangli@umich.edu>
Signed-off-by: Jieming Zhang <jiemingz@nvidia.com>
Signed-off-by: Siddhartha Raman <sraman@nvidia.com>
Signed-off-by: Hollow Man <hollowman@opensuse.org>
Signed-off-by: Phuong Nguyen <phuonguyen@nvidia.com>
Co-authored-by: Vedaanta Agarwalla <142048820+vedaanta@users.noreply.github.com>
Co-authored-by: Hwanseo Choi <hwanseoc@nvidia.com>
Co-authored-by: Vincent <vinnietombari@gmail.com>
Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Co-authored-by: Ziang Li <ziangli@umich.edu>
Co-authored-by: Yang Xu <38851819+YangXu1990uiuc@users.noreply.github.com>
Co-authored-by: Yang Xu <yanxu@nvidia.com>
Co-authored-by: Brandon Zhang <31413216+brandonfzhang@users.noreply.github.com>
Co-authored-by: Jane (Jiancheng) Liu <liujane@nvidia.com>
Co-authored-by: Yanqin Zhai <yanqinz@nvidia.com>
Co-authored-by: Emil Gilliam <egilliam@nvidia.com>
Co-authored-by: Jimmy Zhang <133159885+jiemingz@users.noreply.github.com>
Co-authored-by: jiayus-nvidia <jiayus@nvidia.com>
Co-authored-by: mingyangw <mingyangw@nvidia.com>
Co-authored-by: Takeshi Watanabe <take-cheeze@users.noreply.github.com>
Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com>
Co-authored-by: Mingyang Wang <35635157+saltyminty@users.noreply.github.com>
Co-authored-by: Shraiysh <svaishay@nvidia.com>
Co-authored-by: Jie Fang <jief@nvidia.com>
Co-authored-by: Siddhartha Raman Sundara Raman <sraman@nvidia.com>
Co-authored-by: yeliu-oss <yeliu@nvidia.com>
Co-authored-by: Vincent <34876120+Vinnie6167@users.noreply.github.com>
Co-authored-by: Dimitar (Mitko) Asenov <dimitar.asenov@gmail.com>
Co-authored-by: ℍ𝕠𝕝𝕝𝕠𝕨 𝕄𝕒𝕟 <j88437182@hotmail.com>
Co-authored-by: Josh Park <89948656+jhjpark@users.noreply.github.com>
Co-authored-by: yanzhuo607 <yanzhuoc@nvidia.com>
Co-authored-by: Haisha Zhao <33570593+Hyaloid@users.noreply.github.com>
Co-authored-by: Vince (Junghoon) Han <vincejhan98@gmail.com>
Co-authored-by: cjerry <cjerry@nvidia.com>
Co-authored-by: Zhenyu <370166800@qq.com>
Co-authored-by: Vedaanta Agarwalla <vagarwalla@nvidia.com>
Co-authored-by: Phuong Nguyen <phuonguyen@nvidia.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

cat-infra Build, packaging, tooling, dependency, release, or repository maintenance work. mod-frontend cuDNN frontend APIs, operation graph construction, plans, and user-facing wrappers. orig-qa Reported or owned by quality assurance, validation, or testing.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[Bug]Some broken link target issue found when run document check for cudnn latest code

2 participants