Python-native cudnn.pygraph: graph IR + pluggable execution backends - #336
Merged
Anerudhan merged 38 commits intoJul 7, 2026
Merged
Conversation
…posal)
Modernize the Python-native graph API into the backend-dispatch architecture
from the Frontend v1 "Python API Engine and Graph API Unification" proposal.
Graph construction stays backend-agnostic; a backend is chosen by a first-class
Router at create_execution_plans() time (per Anerudhan's feedback), and the
backend-specific representation (e.g. the C++ cuDNN graph) is generated lazily
only then:
Python Graph API -> create_execution_plans() -> Router -> selected backend
(native engine, else cuDNN)
Layers kept separate:
- Graph IR (Node/Tensor/NativeGraph): engine-agnostic op DAG, full introspection
- BaseEngine: the backend contract (check_support/execute/get_workspace_size +
priority); cuDNN Graph is one routed backend, not a hardcoded default
- Router (engines/router.py): first-supporting by priority; None => cuDNN
Included: the IR, BaseEngine, Router, a CPU-only ReferenceMatmulEngine
(CI-testable correctness oracle), the optional MatmulCuTileEngine, and node
builders for block-scale / MoE / reduction so a DSL fusion backend can consume
them via graph.nodes (replacing the monkey-patch "recorder").
Deferred to follow-ups (see docs/python_native_graph_router.md):
NativeGraph.from_pygraph() (raises NotImplementedError for now), the DSL fusion
backend port, attention backends, and cuDNN lowering of the new node types.
Tests: 42 passing on CPU (IR + Router + reference-engine execute + cuDNN
fallback); cuTile path gated to SM100.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Contributor
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
📝 WalkthroughWalkthroughThis PR adds a Python-native cuDNN graph IR, execution-engine abstractions and routing, a cuTile matmul backend, package exports, and tests for native execution plus cuDNN lowering. ChangesNative graph IR and execution engines
Estimated code review effort: 4 (Complex) | ~75 minutes 🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
…-list routing - NodeType now lists only the op types this version exercises; drop the unused norm/reshape/slice/etc. entries (re-add per-op when needed, following the block-scale / MoE / reduction examples). - Document the target routing model: create_execution_plans() takes one mixed candidate list (native engines + cuDNN heur_modes) and produces a ranked list of plans across backends; this PR ships the first-supporting-by-priority form. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…(unused) Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Drop CONV_FPROP / CONV_DGRAD / CONV_WGRAD: enum entries, the conv_fprop / conv_dgrad builders, their dim inference in nodes.py, cuDNN lowering branches, and the conv test. Re-add per-op when a backend needs conv, following the block-scale / MoE / reduction examples. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Avoid naming specific internal backends in public docs/docstrings. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…NN wrapper) Replace the single-selected-backend + "if native else cpp" fork with the engine-id model: python engines and cuDNN backend engines share one flat id space. Python engines occupy a reserved high region (engine_ids.py, PYTHON_ENGINE_ID_BASE = 1<<20) and each declares a stable engine_id it owns, so ids never shift with registration order (reproducible autotune / pinned plans). - engine_ids.py: PYTHON_ENGINE_ID_BASE + is_python_engine() + a phase-1 CUDNN_HEURISTIC_ENGINE_ID sentinel. Single source of truth for the namespace. - Router.select()->one-engine becomes Router.plan()->ranked list of PlanConfig(engine_id, knobs): supporting python engines (by id) + one trailing cuDNN entry. TODO: interleave the true per-engine cuDNN configs (get_engine_and_knobs_at_index) + real heuristics ranking; for now just concat. - NativeGraph: _selected(engine) -> _plans(list) + _plan_index; add get_execution_plan_count() / select_plan(i). check_support / build_plans / get_workspace_size / execute all dispatch on the selected plan's id via is_python_engine — one predicate, no fork. cuDNN is lowered lazily only when a cuDNN-id plan is selected (pure-python when a python plan wins). - BaseEngine: drop `priority`, add stable `engine_id` (reserved region). reference_matmul = BASE+0, matmul_cutile = BASE+1. Tests updated to assert the plan list; 41 pass on CPU incl. cuDNN fallback. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ront door) Users keep the classic API — g = cudnn.pygraph(...) is unchanged for every existing sample — yet a graph transparently routes to a registered python engine when it's fully represented. No new user-facing class, no rename. pygraph_engines.install(pygraph) (called from __init__, same sanctioned pattern as pygraph.execute = _execute) augments the pybind class in place: - Per-graph mirror (WeakKeyDictionary) records a Node/Tensor IR alongside the real C++ calls for a curated represented set (matmul + common pointwise), mirrored via the NativeGraph builders so the recorded op is exactly what engines consume. - Every other op-builder is auto-wrapped to flag the graph "opaque" — the safe direction: only disables the python path, never changes classic output. - Lifecycle (create_execution_plans/check_support/build_plans/get_workspace_size/ execute/build) routes to a python engine iff one is registered AND the whole graph is represented AND it supports the graph; else delegates to the untouched C++ path. Verified on an L40S against the real cuDNN build: a classic matmul runs byte-identically with and without the augmentation, and a matmul+bias+relu graph built via cudnn.pygraph + ReferenceMatmulEngine routes to the python engine with exact results. Eager for now (C++ graph still built); lazy/pure-python is the follow-up (needs a structured builder per op — multi-tensor returns like sdpa can't be mirrored generically). NativeGraph stays as the standalone/greenfield authoring object sharing the same IR + engines. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Anerudhan
added
cat-feature
…hase 1) Toward the native cudnn.pygraph migration (GEMM-family first). Make the native build->lower->cuDNN execute path actually work end to end, and extend lowering coverage to the GEMM family. Fixes (all latent — the cuDNN execute path had never been GPU-tested): - Thread the cuDNN handle: NativeGraph(handle=...) -> passed to the lowered cudnn.pygraph so heuristics/build have a handle. - Propagate the IR uid to the C++ tensor (was uid=-1 for auto tensors), so execute()'s variant pack (keyed by IR uid) actually binds the buffers. - POINTWISE lowering: the C++ pygraph has no generic pointwise(); dispatch on the mode to the named ops (relu/gelu/sigmoid/tanh, add/mul/sub/div; add/mul also cover bias/scale via broadcast). Lowering coverage added: reduction, block_scale_dequantize, block_scale_quantize (2 outputs), moe_grouped_matmul. Validated on GPU (SM89): matmul and matmul+bias+relu built natively via NativeGraph, lowered to cuDNN, execute with exact parity (new test_native_cudnn_lowering.py, GPU-gated). Full native/router/pygraph suite: 45 passing. Per-op output-shape inference (e.g. reduction reduced dims) and block-scale/moe execution parity are the next slices. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…amily phase 2) - reduction(): take an explicit reduced `dim` (cuDNN requires the reduction output dims set); lowering sets set_dim/set_stride on the cuDNN op. Validated matmul -> reduction(ADD over N) parity on GPU. - lower_tensor(): propagate reordering_type to _make_tensor (e.g. F8_128x4), needed for block-scale scale-factor tensors. Native/router/pygraph + GPU parity suite: 46 passing. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…s (phase 3) Complete the GEMM-family native lowering with block-scale, validated on SM100. Two more latent cuDNN-path bugs fixed: - _lower_to_cpp passed io_data_type=None -> cudnn.pygraph rejects None. Now omit io when unset; default intermediate/compute to FLOAT (matching cudnn.graph()) so cuDNN infers virtual (intermediate) tensor dtypes during build. - lower_tensor now propagates reordering_type (F8_128x4) and omits data_type when unset (NOT_SET) so cuDNN infers fused block-scale dequant output types. Validated on SM100: dequant(A_fp4)@dequant(B_fp4) with F8_128x4 SFs builds + executes via NativeGraph (test gated to SM100 + torch fp4; parity harness = the repo's own fp4 test, which also only checks execution). CPU overhead of the native Python layer (512^3 fp16, L40S): build +0.40 ms on ~106 ms (~0.4%, dominated by cuDNN heuristics); execute +0.3 us/call (9.8 -> 10.1 us). Negligible. Native/router/pygraph + GPU parity (matmul, bias+relu, reduction, block-scale): 48 passing. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…y complete) - Add moe output-shape inference (token [1,T,H], weight [E,H,N] -> out [1,T,N]) so NativeGraph.validate() passes; cuDNN infers the same at build. - GPU parity test (self-contained per-expert reference; no dependency on the upstream test's helper) — validated on SM100. GEMM family now fully native-lowered + validated on GPU: matmul, pointwise (bias/relu), reduction, block-scale nvfp4, moe. Suite: 48 passing. Next: non-GEMM ops (norms/reshape/slice/...) then the C++ _op rename + atomic flip of cudnn.pygraph. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
… (first norm) Systemic fix: op-created C++ tensors (op outputs / virtuals) get uids assigned by the C++ FE during build_operation_graph, in ITS enumeration order — which does not match IR allocation order for multi-output ops (rmsnorm assigns INV_VARIANCE=5, Y=6 while the IR allocated Y=5, inv_var=6). Keying the variant pack by raw IR uids bound Y's buffer to inv_var: a [N,C,H,W] fp16 write into a 16-byte buffer (heap corruption / NaN). Single-output ops only worked by allocation-order coincidence. Fix: keep the lowering tensor_map; after build_operation_graph query every C++ tensor's real uid into an explicit IR-uid -> C++-uid map; execute() translates variant-pack keys through it. No more order coincidence anywhere. rmsnorm added as the first-class norm template (per "no corner-cutting" — the generic opaque-op bridge was rejected/reverted since it makes non-GEMM ops un-introspectable black boxes): named input/scale/epsilon/bias ports, Y/inv_var outputs, norm_forward_phase param, pass-by-value epsilon; Y/inv_var dims carried in the IR, cuDNN infers on its side. GPU parity: errY=0.0019, errI=0.0. Suite: 49 passing (GEMM family re-validated through the translation path). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Systematic uid review — four assignment paths existed:
1. user at creation: tensor(uid=...) (pybind _make_tensor, default -1)
2. user post-creation: tensor.set_uid() (mainline integrator pattern)
3. C++ FE auto-assign at build_operation_graph (enumeration order,
nondeterministic for multi-output ops) <- the coincidence trap
4. Python IR _alloc_uid (eager, sequential)
New invariant: for Python-built graphs, (3) NEVER triggers. The IR assigns
every uid eagerly at creation (auto or user-specified); lowering pushes ALL of
them explicitly to C++ — inputs via _make_tensor(uid=), op-created
outputs/virtuals via one set_uid loop over the complete tensor_map (single
point, impossible to forget per-op). Mixed construction (extending the lowered
C++ graph directly) is unsupported: a graph is pure-Python or pure-C++.
- Replace the IR->C++ uid translation map with a post-build ASSERTION: a
lowering path that fails to push a uid now fails loudly instead of being
silently translated (or worse, mis-binding buffers).
- _alloc_uid skips user-reserved uids; duplicate explicit uids rejected eagerly
at tensor() (C++ would only fail at build).
- execute() keys the variant pack by IR uids directly (== C++ uids by
construction).
Suite: 50 passing on SM100 (rmsnorm multi-output canary + block-scale included).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
… method name Cover the entire pointwise surface of the C++ pygraph (54 methods) natively: - Canonical op kind: params["mode"] IS the C++ pygraph method name (the pointwise_mode enum is not exposed to Python; the method name is the semantic name). Lowering collapses to a direct getattr dispatch — the mode<->method mapping table is deleted as a concept. - 47 uniform ops are generated from _POINTWISE_TENSOR_ARGS, a table of the pybind tensor-argument names per op (mirrors the C++ signatures), so both positional and the classic keyword call styles (bias(input=, bias=), max(input0=, input1=)) work — required for the eventual cudnn.pygraph flip. - 7 ops with scalar attributes get explicit builders storing them in params (introspectable): relu(negative_slope/lower_clip/upper_clip), leaky_relu, swish(swish_beta), gen_index(axis), + relu/leaky_relu/swish backwards. Lowering forwards them as keywords. - ReferenceMatmulEngine: keys move to method names; declines pointwise nodes carrying scalar attributes it does not implement (correct-by-construction). - Front-door mirror: classic calls passing scalar extras (e.g. relu clips) now flag the graph opaque instead of silently dropping the attribute and mis-routing to a python engine. Tests: every builder exercised in both call styles + scalar-attr introspection (CPU); sqrt/abs/max/min chain through real cuDNN on GPU. 53 passing. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…lowering) All norms native — rmsnorm(_backward), layernorm(_backward), adalayernorm(_backward), instancenorm(_backward), batchnorm, batchnorm_inference, batchnorm_backward — through ONE mechanism instead of per-op code: - _STRUCTURED_OPS: a declarative table per op — NodeType, tensor-input ports (== the C++ pybind kwarg names), enum/scalar params (norm_forward_phase, has_dbias), output ports in C++ return order, and per-output shape inference (IR-side dims for introspection; cuDNN re-infers at build). Builders are generated (keyword call style, as these ops are used repo-wide); lowering is one generic branch: kwargs assembly + one call + zip outputs. - List inputs (batchnorm peer_stats) become indexed ports (peer_stats_i) + a count param, reassembled at lowering. - The hand-written rmsnorm builder AND its lowering branch are deleted — migrated into the table; the suite re-validates rmsnorm through the generic path (multi-output uid canary intact). GPU parity: layernorm fwd (Y/mean/inv_var) + layernorm_backward (DX/DScale/ DBias) vs torch autograd, using the supported LN config ([N,C,1,1] channels_last, as in classic test_layernorm — the initial row-major 4D attempt fails identically on the classic API, i.e. a kernel-support limit, not a lowering bug). CPU: every table op builds a first-class node with named ports; peer_stats port machinery covered. 56 passing. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
… one table
_STRUCTURED_OPS now covers 25 ops — norms (11 incl. genstats), reduction,
block-scale (de)quantize, moe fwd/bwd, conv fprop/dgrad/wgrad, reshape, slice,
transpose, concatenate, rope fwd/bwd — one declarative entry each, one generic
lowering branch. Only matmul (positional ergonomics + front-door mirror) and
sdpa fwd/bwd (conditional kwarg assembly) remain explicit.
Deleted in the collapse: the hand-written reduction / block_scale_dequantize /
block_scale_quantize / moe_grouped_matmul builders AND their four lowering
branches, plus nodes.py moe shape inference (moved to the table). The suite
re-validates all of them through the generic path on GPU.
Table mechanics extended (each a one-word spec key, no new concepts):
- attrs: scalar/enum/list params forwarded verbatim (padding vectors, axis,
slices, permutation, reshape_mode, rope_dim, mode, ...). Conv accepts BOTH
the symmetric `padding` convenience and pre/post_padding — forwarded as
given; pybind overload resolution picks the right C++ binding.
- out_dims reserved kwarg (list, or {port: dims}): explicit output shapes for
ops cuDNN cannot infer — generalizes reduction's old `dim` param.
- push_output_dims: IR dims pushed to C++ for dgrad/wgrad/reduction/reshape/
moe_bwd (classic API also requires set_dim there).
- no_cdt: bindings without compute_data_type (reshape, concatenate).
- Builders accept tensors positionally or by port name; infer lambdas are
best-effort (try/except -> None; C++ validates at build).
GPU parity added: conv_fprop vs torch conv2d (NHWC), incl. asserting the
table's shape inference. CPU: all 25 ops x 2 call styles + out_dims. 58 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…urface The six sdpa variants (sdpa, sdpa_backward, sdpa_fp8, sdpa_fp8_backward, sdpa_mxfp8, sdpa_mxfp8_backward) are now declared in _CAPTURED_OPS, the third and final table mechanism: builders capture ALL kwargs generically — tensor values (incl. torch/dlpack) become named ports (port == C++ kwarg), scalars / enums / score_mod callbacks go to params verbatim, dropout tuples are flattened per element — and lowering rebuilds the kwargs for one C++ call. The full C++ kwarg surface (~130 args: paged attention tables, diagonal bands, sink tokens, cu_seqlens, fp8 descales/amaxes, ...) is supported without hand-mirroring any of it, and future binding args are picked up automatically. Deleted: the explicit sdpa/sdpa_backward builders (~170 lines, common-args only) + their two lowering branches + nodes.py sdpa shape inference (moved to table lambdas — and fixed: O is q-shaped with v's head dim, not v-shaped). Semantics now match the classic API exactly: sdpa always returns (O, Stats) with Stats None in inference mode (generate_stats/is_inference logic); output dim/stride are pushed to C++ (the SDPA node requires O's layout pre-validate — that's how BSHD vs BHSD output is chosen). GPU: sdpa causal fp16 EXECUTION parity vs torch SDPA (was build-only before). 59 passing. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The public cudnn.pygraph name now binds the Python IR class (class name:
pygraph; module: python/cudnn/pygraph.py — no "relative-to-history" naming).
The C++ graph builder is internal-only at cudnn._pybind_module.pygraph and is
reached exclusively through lowering: a graph is pure-Python or pure-C++,
never mixed. Zero C++ changes — the demotion is by namespace, not rebuild.
Deleted in the flip (afterthought residue):
- pygraph_engines.py front-door + its tests (no install()/monkey-patching
anywhere: register_backend is a native method on the class)
- NativeGraph.from_pygraph stub (meaningless now), use_native back-door
- docs/python_native_graph_router.md (initial-brainstorm doc, per review)
Drop-in surface for classic parity, driven by iterating the repo's own test
files until green (each item below was a real failure caught and fixed):
- conditional outputs ("maybe"): rmsnorm_backward(has_dbias=False) -> DBias
None; norm fwd INFERENCE -> mean/inv_var None; batchnorm next_running_*
present iff in_running_* given (classic returns None for absent outputs)
- torch interop: tensor(dim=x.size()) (torch.Size), data_type=torch.bfloat16
(converted at the C++ boundary via _library_type, IR stores user's value)
- output dtype semantics: an output without explicit set_data_type gets io
dtype (was mis-defaulted to intermediate FLOAT -> fp32 into fp16 buffers)
- Tensor gains the classic setter/getter surface (set_ragged_offset,
set_reordering_type, set_is_pass_by_value, ...); tensor_like(cudnn tensor);
tensor_scalar; CPU tensor_like -> pass-by-value (classic rule)
- ragged (THD) output layout: outputs' ragged_offset now pushed to C++ at all
mapping sites (was silently dense -> wrong values in sdpa_thd)
- validate-time table shape inference (topological): chained ops whose inputs
are virtual (conv on a relu output) infer once inputs are known;
builder-time infer stays as best-effort for direct inputs
- classic lifecycle: build_operation_graph lowers eagerly when no python
engines are registered, so deselect_*/query methods work between classic
steps via __getattr__ delegation to the lowered graph; build_plans(policy)
passthrough; deserialize(*args, **kwargs) passthrough incl.
enforce_precompiled; execute override_uids/shapes/strides + dlpack pointers;
get_execution_plan_count = python engines + backend's dynamically-queried
count (frontend NEVER statically enumerates backend engines — they vary by
backend version; Router keeps ONE delegating cuDNN entry by design)
- stride optional after set_dim (row-major inferred), None variant-pack keys
tolerated, C++-tensor keys resolved via get_uid
Validated: our suite (56) + classic spot-runs all green on real GPUs —
matmul_bias_relu, rmsnorm, layernorm, batchnorm, conv_fprop (incl.
execute_plan_at_index), apply_rope, kernel_cache, sdpa_with_caching, sdpa_thd,
sdpa_chunked_prefill (ragged+paged), conv_genstats, conv_reduction, slice,
block_scale_quantize_dynamic_shape, wgrads. Full-suite runs on SM100 + mhas in
flight; residuals to follow. Known pre-existing env skew (fails identically on
the unflipped installed package): test_deviceless_aot_compilation on this box.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Two classic-parity fixes surfaced by the full mhas run (3567 uniform failures, one root cause): - cudnnGraphNotSupportedError must fire at graph.validate(): the classic test waiver pattern is try/except-skip AROUND validate(), with build_operation_graph() called bare. With no python engines registered, validate() now lowers and runs the C++ validate right there (unsupported configs skip, not fail); build_operation_graph()/plan creation are staged behind flags so each C++ step runs exactly once in classic sequencing. Python-engine graphs still never touch C++ at validate. - compute_data_type=None is now OMITTED at every lowering site (matmul / pointwise / structured / captured) instead of passed through: classic ops default to NOT_SET in C++; pybind rejects None. Also converts via _library_type when set (torch dtype parity). Previously-failing mhas case now skips as on classic; our suite 56 passing. Full-suite + full-mhas reruns in flight. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Ranking policy is intentionally undecided; what IS decided: policy pluggable at three levels (Router subclass / per-graph / process default); plan() may return any ordering or mix; backend engine sets are discovered per graph at plan time (never statically enumerated); PlanConfig can carry concrete backend engine configs, with pygraph._lower_cudnn_plan as the designated point to honor them. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
YangXu1990uiuc
marked this pull request as ready for review
July 2, 2026 21:14
YangXu1990uiuc
changed the title
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Actionable comments posted: 8
🧹 Nitpick comments (7)
test/python/test_graph_native.py (2)
383-398: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueUnused
Obinding.
Ois unpacked at Line 390 but never referenced afterward (onlystatsis checked). Ruff also flagsOas an ambiguous name (E741) alongside similar occurrences at Lines 373 and 487.🧹 Proposed fix
- O, stats = g.sdpa(Q, K, V, is_inference=False, attn_scale=0.125, name="attn") + _, stats = g.sdpa(Q, K, V, is_inference=False, attn_scale=0.125, name="attn")🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@test/python/test_graph_native.py` around lines 383 - 398, The test currently binds the SDPA output to O even though only stats is asserted, and Ruff also treats O as an ambiguous name; update test_sdpa_training in NativeGraph tests to avoid the unused/ambiguous binding by unpacking only the needed value or renaming the placeholder consistently with the similar SDPA tests in this file. Keep the assertions on g.nodes[0].outputs and stats.dim unchanged.Source: Linters/SAST tools
404-415: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueRedundant exception type in
except (ImportError, Exception).
ImportErroris a subclass ofException, so listing both is redundant;except Exceptionalone is equivalent. Alsoerrat Lines 410-411 is unpacked but unused twice.🧹 Proposed fix
- err, device_id = cudart.cudaGetDevice() - err, props = cudart.cudaGetDeviceProperties(device_id) + _, device_id = cudart.cudaGetDevice() + _, props = cudart.cudaGetDeviceProperties(device_id) return props.major * 10 + props.minor >= 100 - except (ImportError, Exception): + except Exception: return False🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@test/python/test_graph_native.py` around lines 404 - 415, The `cutile_available` fixture in `test_graph_native.py` has redundant exception handling because `ImportError` is already covered by `Exception`; simplify the `except (ImportError, Exception)` to just `except Exception`. While there, clean up the `cudart.cudaGetDevice()` and `cudart.cudaGetDeviceProperties()` unpacking so the unused `err` values are not assigned, keeping the fixture focused on checking CUDA tile availability.Source: Linters/SAST tools
test/python/test_native_cudnn_lowering.py (2)
124-156: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueMinor: prefer unpacking over list concatenation.
Ruff RUF005 suggests
[*fto, T]instead offto + [T]at Line 151 — purely stylistic.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@test/python/test_native_cudnn_lowering.py` around lines 124 - 156, The test_native_moe_grouped_matmul_lowers_to_cudnn test uses list concatenation to append T to fto; replace the bounds construction in test_native_cudnn_lowering.py within test_native_moe_grouped_matmul_lowers_to_cudnn with unpacking syntax instead of fto + [T]. Keep the rest of the reference logic unchanged, and update the local bounds variable accordingly so the style aligns with Ruff RUF005.Source: Linters/SAST tools
159-183: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueAmbiguous variable name
O.Flagged by Ruff (E741) at Line 173; consistent with the same pattern in
test_graph_native.py. Domain convention (attention output) makes this understandable, but consider renaming (e.g.Out) if lint compliance matters for CI.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@test/python/test_native_cudnn_lowering.py` around lines 159 - 183, The attention output tensor name `O` in `test_native_sdpa_fwd_lowers_to_cudnn` is flagged as an ambiguous variable name by Ruff, so rename it to a clearer symbol such as `Out` and update the subsequent output setup and execution calls that reference it. Keep the same `g.sdpa(...)` result handling and `set_output` usage, just replace the ambiguous identifier consistently throughout the test.Source: Linters/SAST tools
python/cudnn/engines/matmul_cutile_engine.py (1)
151-158: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winDocstring
Raisessection omitsValueError.
check_supportalso raisesValueErrorfor non-row-major layouts (line 184), but the docstring's Raises block only listsRuntimeErrorandNotImplementedError.📝 Proposed fix
Raises: RuntimeError: If GPU or driver doesn't meet requirements NotImplementedError: If graph contains unsupported operations + ValueError: If a tensor's layout is not row-major contiguous """As per path instructions, "Focus on documentation" for
python/cudnn/**files.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@python/cudnn/engines/matmul_cutile_engine.py` around lines 151 - 158, Update the docstring for check_support in the matmul_cutile_engine module so its Raises section includes ValueError alongside RuntimeError and NotImplementedError. Keep the documentation aligned with the method’s actual behavior, especially the non-row-major layout validation in check_support.Source: Path instructions
python/cudnn/engines/__init__.py (1)
30-35: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick winCatch
ImportErrorspecifically instead of bareException.The internal
try/except ImportErrorguards inmatmul_cutile_engine.pyalready turn missingcuda.tile/cuda.bindingsintoct = None/cudart = Noneat import time — the constructor is what raisesImportErrorfor missing deps. Catching broadExceptionhere means any other error while importing the module (e.g. a real bug introduced later) is silently swallowed too, disabling the engine with no diagnostic.♻️ Proposed fix
try: from .matmul_cutile_engine import MatmulCuTileEngine # noqa: F401 __all__.append("MatmulCuTileEngine") -except Exception: # noqa: BLE001 +except ImportError: MatmulCuTileEngine = None # type: ignore🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@python/cudnn/engines/__init__.py` around lines 30 - 35, The import guard in MatmulCuTileEngine is too broad and is swallowing real module bugs by catching Exception. Update the guarded import in the __init__ module to catch ImportError only, so missing dependencies still disable the engine while unexpected import-time failures surface normally. Keep the MatmulCuTileEngine symbol handling and __all__ update tied to that specific import path.pyproject.toml (1)
79-82: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winConsider pinning
cuda-tileincutile. It is still under active development, so unbounded upgrades can breakMatmulCuTileEngineon a future release.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@pyproject.toml` around lines 79 - 82, The cutile dependency list is using an unpinned cuda-tile requirement, which allows future breaking upgrades. Update the dependency entry for cutile in the dependency configuration to pin cuda-tile to a compatible version, keeping the existing package grouping intact so MatmulCuTileEngine continues to work with known-good releases.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@python/cudnn/engines/base.py`:
- Around line 52-54: The default engine_id in BaseEngine can silently collide
with ReferenceMatmulEngine, so update BaseEngine to require subclasses to
provide a real id instead of inheriting the sentinel. Add a guard in BaseEngine
registration/initialization or make engine_id abstract so any subclass that
leaves PYTHON_ENGINE_ID_BASE unchanged fails fast. Use BaseEngine and
ReferenceMatmulEngine as the key symbols when implementing the validation.
In `@python/cudnn/engines/matmul_cutile_engine.py`:
- Around line 143-149: The MatmulCuTileEngine stores a device in __init__ but
never uses it to verify tensor placement, so update check_support and/or execute
to validate that input and output tensors are on the same CUDA device/context as
self.device before launching kernels. Use the MatmulCuTileEngine, check_support,
and execute paths to compare the tensors’ device against the configured device
and fail fast with a clear error when they do not match.
- Around line 159-170: In MatmulCuTileEngine.__init__ (or the device/driver
validation block), the CUDA runtime return codes from cudaGetDevice,
cudaGetDeviceProperties, and cudaDriverGetVersion are ignored. Check each err
immediately after the call and raise an error with the actual CUDA failure
before using device_id, props, or driver_version. Keep the existing SM100+ and
driver-version validation, but only after confirming each CUDA API call
succeeded so the RuntimeError reflects the real failure source.
In `@python/cudnn/graph_types.py`:
- Around line 76-192: Tensor hashing/equality in the Tensor class is unsafe
because __hash__ and __eq__ rely on the mutable uid field, which defaults to 0
and can change via set_uid. Update the Tensor comparison logic so objects used
in sets/dicts remain stable before and after UID assignment, ideally by making
equality/hash identity-based or by otherwise ensuring uid is immutable once
hashed. Review the Tensor methods set_uid, __hash__, and __eq__ together so the
fix is consistent and prevents collisions between newly created tensors.
In `@python/cudnn/nodes.py`:
- Around line 115-134: The pointwise broadcast inference in _infer_pointwise
currently overwrites accumulated dimensions when it sees a higher-rank input and
never validates same-rank broadcast compatibility. Update the logic in
_infer_pointwise so it merges all input shapes progressively instead of
replacing max_dim, aligning ranks as needed and combining each dimension with
broadcast rules (equal or 1), and raise or handle incompatible dimensions rather
than silently dropping previously inferred sizes.
- Around line 82-113: The batch-dimension handling in _infer_matmul is silently
merging incompatible shapes by taking max(a_val, b_val) instead of validating
broadcast compatibility. Update _infer_matmul (and, if needed, _validate_matmul)
so batch dims are checked per position: allow equal dims or 1-to-other
broadcasting, and raise a shape error for mismatches like 2 vs 3 rather than
inferring c.dim from max.
In `@python/cudnn/pygraph.py`:
- Around line 1-17: The new pygraph/NativeGraph public API is missing both test
coverage and the user-facing router documentation. Add fe_api tests under
test/python/fe_api that exercise the pygraph and NativeGraph entry points and
their routing behavior, and create docs/python_native_graph_router.md to
describe how create_execution_plans(), Router, and backend selection work. Use
the symbols pygraph, NativeGraph, create_execution_plans(), and Router to align
the tests and docs with the new flow.
- Around line 212-227: Make `scalar_type` required in `tensor_scalar` so
pass-by-value scalars cannot be created without an explicit type and lose their
embedded value during lowering. Update the `PyGraph.tensor_scalar` method
signature to remove the default for `scalar_type`, and keep the `Tensor(...,
is_pass_by_value=True, pass_by_value=value, scalar_type=scalar_type)`
construction aligned with that contract. If you prefer to support omission
instead, ensure the omitted path still preserves `pass_by_value`, and add a
`test/python/fe_api` case covering the no-`scalar_type` call site.
---
Nitpick comments:
In `@pyproject.toml`:
- Around line 79-82: The cutile dependency list is using an unpinned cuda-tile
requirement, which allows future breaking upgrades. Update the dependency entry
for cutile in the dependency configuration to pin cuda-tile to a compatible
version, keeping the existing package grouping intact so MatmulCuTileEngine
continues to work with known-good releases.
In `@python/cudnn/engines/__init__.py`:
- Around line 30-35: The import guard in MatmulCuTileEngine is too broad and is
swallowing real module bugs by catching Exception. Update the guarded import in
the __init__ module to catch ImportError only, so missing dependencies still
disable the engine while unexpected import-time failures surface normally. Keep
the MatmulCuTileEngine symbol handling and __all__ update tied to that specific
import path.
In `@python/cudnn/engines/matmul_cutile_engine.py`:
- Around line 151-158: Update the docstring for check_support in the
matmul_cutile_engine module so its Raises section includes ValueError alongside
RuntimeError and NotImplementedError. Keep the documentation aligned with the
method’s actual behavior, especially the non-row-major layout validation in
check_support.
In `@test/python/test_graph_native.py`:
- Around line 383-398: The test currently binds the SDPA output to O even though
only stats is asserted, and Ruff also treats O as an ambiguous name; update
test_sdpa_training in NativeGraph tests to avoid the unused/ambiguous binding by
unpacking only the needed value or renaming the placeholder consistently with
the similar SDPA tests in this file. Keep the assertions on g.nodes[0].outputs
and stats.dim unchanged.
- Around line 404-415: The `cutile_available` fixture in `test_graph_native.py`
has redundant exception handling because `ImportError` is already covered by
`Exception`; simplify the `except (ImportError, Exception)` to just `except
Exception`. While there, clean up the `cudart.cudaGetDevice()` and
`cudart.cudaGetDeviceProperties()` unpacking so the unused `err` values are not
assigned, keeping the fixture focused on checking CUDA tile availability.
In `@test/python/test_native_cudnn_lowering.py`:
- Around line 124-156: The test_native_moe_grouped_matmul_lowers_to_cudnn test
uses list concatenation to append T to fto; replace the bounds construction in
test_native_cudnn_lowering.py within
test_native_moe_grouped_matmul_lowers_to_cudnn with unpacking syntax instead of
fto + [T]. Keep the rest of the reference logic unchanged, and update the local
bounds variable accordingly so the style aligns with Ruff RUF005.
- Around line 159-183: The attention output tensor name `O` in
`test_native_sdpa_fwd_lowers_to_cudnn` is flagged as an ambiguous variable name
by Ruff, so rename it to a clearer symbol such as `Out` and update the
subsequent output setup and execution calls that reference it. Keep the same
`g.sdpa(...)` result handling and `set_output` usage, just replace the ambiguous
identifier consistently throughout the test.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: b587980a-b2e5-4383-a997-eaa5f969bf73
📒 Files selected for processing (14)
pyproject.tomlpython/cudnn/__init__.pypython/cudnn/engines/__init__.pypython/cudnn/engines/base.pypython/cudnn/engines/engine_ids.pypython/cudnn/engines/matmul_cutile_engine.pypython/cudnn/engines/reference_matmul_engine.pypython/cudnn/engines/router.pypython/cudnn/graph_types.pypython/cudnn/nodes.pypython/cudnn/pygraph.pytest/python/test_engine_router.pytest/python/test_graph_native.pytest/python/test_native_cudnn_lowering.py
…ew items 2, 6) Review item 2 (reproduced bugs): - ONE plan index space: [0, n_python) are python plans, [n_python, ...) are the backend's plans (sub-index = index - n_python, queried dynamically). get_execution_plan_count() and select_plan() now agree; selecting a backend sub-index lowers on demand, builds via build_plan_at_index and executes via _execute_plan_at_index (sub-index 0 == the classic default path). - select_plan() survives build()/execute(): build() no longer silently re-plans when a plan list exists (explicit create_execution_plans() still re-plans). Review item 6: - register_backend() validates at registration: engine_id must be a stable int in the reserved python region, unique per graph; registration after planning is rejected. BaseEngine.engine_id defaults to None so a subclass that forgets to declare identity fails clearly instead of silently colliding. - Decline signal narrowed: an engine declines ONLY via NotImplementedError or cudnn.cudnnGraphNotSupportedError (the classic unsupported-graph signal); ValueError/RuntimeError now propagate as engine bugs instead of silently falling back to cuDNN. Reference/cuTile engines updated accordingly. Regression tests for all of the above (pin-survives-execute, duplicate/missing id, post-planning registration, unexpected-exception propagation). 59 passing + classic spot files green. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ew item 1) The engine contract now represents a real JIT/DSL backend: - propose_plans(graph) -> [PlanConfig]: one engine may expose several configurations to ranking/autotune (default: one plan with default_knobs when check_support accepts). PlanConfig moves to engines/base.py. - build_plan(graph, plan) -> CompiledPlan: the expensive JIT step, run ONCE per (graph, selected plan) at build_plans() time. The compiled artifact is cached ON THE GRAPH (keyed by plan index), so one engine instance is safely reusable across graphs and repeated execution reuses the artifact. The selected plan's knobs reach build_plan verbatim. - CompiledPlan.get_workspace_size(): plan-specific workspace; graph get_workspace_size() reports it for python plans. - ExecutionContext(handle, stream, workspace, override_uids/shapes/strides) passed to CompiledPlan.execute(): stream resolved from the caller's handle (classic cudnn.set_stream semantics); caller workspace object reaches the plan; no engine hard-codes a stream (cuTile now launches on ctx.stream). - Simple eager engines are unchanged in spirit: implement execute() only; the default build_plan wraps it in a trivial CompiledPlan. Acceptance tests per the review: two knob proposals from one engine with the selected plan's knobs observed at build+execute; compile-once artifact reuse across executions; same engine instance on two graphs without state collision; plan-specific nonzero workspace; caller workspace object identity at execute. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Contributor
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (2)
python/cudnn/engines/matmul_cutile_engine.py (2)
143-149: 📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick winInstall instructions in the class docstring are inconsistent with the
ImportErrormessage.The docstring (Line 137) tells users to
pip install cuda-tile, but theImportErroron Line 146 instructspip install nvidia-cudnn-frontend[cutile]. These should match to avoid confusing users about the correct way to enable this engine.📝 Proposed doc alignment
Requirements: - Blackwell GPU (SM100+) - CUDA Toolkit 13.1+ - - cuda-tile package: pip install cuda-tile + - cuda-tile package: pip install nvidia-cudnn-frontend[cutile] """As per path instructions ("Focus on documentation").
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@python/cudnn/engines/matmul_cutile_engine.py` around lines 143 - 149, The class docstring for MatmulCuTileEngine and its __init__ ImportError message give different install commands, so align them to the same package name and usage. Update the documentation near MatmulCuTileEngine so the install instructions match the ImportError text (or vice versa), keeping the guidance consistent for both the cuda-tile and cuda-python requirements.Source: Path instructions
151-219: 📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick winAdd a gated
test/python/fe_apicase forMatmulCuTileEngineThat tree has no test that exercises this engine or its Blackwell-onlycheck_support/executepath.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@python/cudnn/engines/matmul_cutile_engine.py` around lines 151 - 219, Add a gated test case under test/python/fe_api that actually exercises MatmulCuTileEngine, since nothing there currently covers its Blackwell-only path. Use MatmulCuTileEngine.check_support and MatmulCuTileEngine.execute in the test to verify the support gate and a basic execution flow, with the test conditionally skipped or marked for SM100+/r580+ environments so it only runs where the engine is supported.Source: Path instructions
🧹 Nitpick comments (1)
python/cudnn/engines/matmul_cutile_engine.py (1)
202-219: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick winAvoid
assertfor runtime data validation.The inner-dimension and batch-size checks (Lines 205, 213, 214) validate runtime tensor shapes from
tensor_data, butassertstatements are stripped when Python runs with-O. In an optimized build a mismatchedK/batchwould slip through and launch a kernel against incompatible buffers instead of failing fast. Prefer explicitValueErrorraises to match theelsebranch on Line 219.♻️ Proposed change
if a.ndim == 2: M, K = a.shape K2, N = b.shape - assert K == K2, f"Inner dimensions must match: {K} vs {K2}" + if K != K2: + raise ValueError(f"Inner dimensions must match: {K} vs {K2}") grid = (ct.cdiv(M, TM), ct.cdiv(N, TN), 1) ct.launch(stream, grid, _get_matmul_kernel(), (a, b, c, M, N, K, TM, TN, TK)) elif a.ndim == 3: batch, M, K = a.shape batch2, K2, N = b.shape - assert batch == batch2, f"Batch sizes must match: {batch} vs {batch2}" - assert K == K2, f"Inner dimensions must match: {K} vs {K2}" + if batch != batch2: + raise ValueError(f"Batch sizes must match: {batch} vs {batch2}") + if K != K2: + raise ValueError(f"Inner dimensions must match: {K} vs {K2}")🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@python/cudnn/engines/matmul_cutile_engine.py` around lines 202 - 219, In the matmul launch path inside the engine method that handles a.ndim == 2 and a.ndim == 3, replace the runtime shape checks currently using assert with explicit ValueError raises so they are not removed under optimized Python. Update the inner-dimension validation for K/K2 and the batch-size validation for batch/batch2 to fail fast with clear messages, matching the existing Unsupported tensor dimensions handling in the same function.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Outside diff comments:
In `@python/cudnn/engines/matmul_cutile_engine.py`:
- Around line 143-149: The class docstring for MatmulCuTileEngine and its
__init__ ImportError message give different install commands, so align them to
the same package name and usage. Update the documentation near
MatmulCuTileEngine so the install instructions match the ImportError text (or
vice versa), keeping the guidance consistent for both the cuda-tile and
cuda-python requirements.
- Around line 151-219: Add a gated test case under test/python/fe_api that
actually exercises MatmulCuTileEngine, since nothing there currently covers its
Blackwell-only path. Use MatmulCuTileEngine.check_support and
MatmulCuTileEngine.execute in the test to verify the support gate and a basic
execution flow, with the test conditionally skipped or marked for SM100+/r580+
environments so it only runs where the engine is supported.
---
Nitpick comments:
In `@python/cudnn/engines/matmul_cutile_engine.py`:
- Around line 202-219: In the matmul launch path inside the engine method that
handles a.ndim == 2 and a.ndim == 3, replace the runtime shape checks currently
using assert with explicit ValueError raises so they are not removed under
optimized Python. Update the inner-dimension validation for K/K2 and the
batch-size validation for batch/batch2 to fail fast with clear messages,
matching the existing Unsupported tensor dimensions handling in the same
function.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 077280ea-d4a3-4f90-9aab-cd118e392076
📒 Files selected for processing (7)
python/cudnn/engines/__init__.pypython/cudnn/engines/base.pypython/cudnn/engines/matmul_cutile_engine.pypython/cudnn/engines/reference_matmul_engine.pypython/cudnn/engines/router.pypython/cudnn/pygraph.pytest/python/test_engine_router.py
🚧 Files skipped from review as they are similar to previous changes (2)
- python/cudnn/engines/init.py
- python/cudnn/engines/reference_matmul_engine.py
…ides are cuDNN The frontend Python graph is as much cuDNN as the C++ library; identifiers that used 'cudnn' to designate the backend side now say 'backend': - CUDNN_HEURISTIC_ENGINE_ID -> BACKEND_HEURISTIC_ENGINE_ID - _lower_cudnn_plan / _has_cudnn_plan / _cudnn_heuristics -> _lower_backend_plan / _has_backend_plan / _backend_heuristics - _assert_ran_on_cudnn -> _assert_ran_on_backend - test_native_cudnn_lowering.py -> test_native_backend_lowering.py (tests *_lowers_to_cudnn -> *_lowers_to_backend, mixed-router / one-shot test names likewise) - docstrings/comments: 'cuDNN entry/sentinel/slot/path/side' -> 'backend ...' throughout; 'the cuDNN C++ backend' stays where it describes what the backend is. Also fixes a stale TYPE_CHECKING import left by the module rename. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Anerudhan
approved these changes
Jul 6, 2026
Collaborator
|
@cudnn-ci-bot run |
|
🚀 Running mirror pipeline Branch: cudnn-gh/pr-336-921dcf5 |
…per, labels, naming, layout truth Root-caused from the internal CI failures (py_samples / pycudnnTest); every item below reproduces 1:1 against the classic package on the same GPU/backend and is fixed + validated (pycudnnTest 26/26, all 13 CI sample notebooks pass, local battery 2172/0): - Constructor and tensor() are POSITIONALLY IDENTICAL to the classic API (name is the constructor's first positional arg — pycudnnTest passes it positionally; classic sm_count/sm_version/kernel_cache/device_property/ dynamic-shape params explicit; classic tensor() order with is_pass_by_value/ ragged_offset/reordering before name/uid; NOT_SET/-1/NONE sentinels normalized). New params (backends/router) are keyword-only. Guarded by test_api_signature_parity, which reads the classic order from the pybind docstring/wrapper itself. - wrapper.py (cudnn.Graph) recognizes IR tensors: one _GRAPH_TENSOR_TYPES tuple replaces 7 isinstance(cudnn.tensor) sites (the notebooks' silent UnboundLocalError/mis-capture). - Duplicate tensor names are legal classic LABELS (pycudnnTest builds two 'weight's): uid is identity; the name index serves unique names only and ambiguous-name lookups raise instead of guessing. - Op outputs are auto-named with the classic C++ conventions (node::MEAN/INV_VARIANCE/DSCALE..., per-op overrides for rmsnorm_backward's ::Dscale/::Dbias) — wrapper.Graph canonical-name lookups depend on them. - Multi-output ops return a LIST like classic pybind (pycudnnTest dispatches on isinstance(res, list)). - Layout truth: backend-inferred dim/stride are reflected back into the IR after build_operation_graph (_sync_ir_shapes_from_backend) — wrapper allocates output buffers from IR getters; provisional row-major strides are no longer observable post-build. push_output_dims ops push stride only when USER-assigned (pushing inferred row-major into an NHWC graph made the backend reject dgrad+add fusion). - tensor_like normalizes non-torch DLPack objects (CuPy .strides is in BYTES) through torch.from_dlpack — NHWC CuPy inputs no longer silently become row-major. - get_data_type() returns the cudnn enum when the user stored a torch dtype (classic converts at set time). - validate() no longer auto-marks leaf outputs as non-virtual — discarding a result (training SDPA's Stats in the paged sample) is legal classic usage; auto-marking made its uid required in the variant pack. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Collaborator
|
@cudnn-ci-bot run |
|
🚀 Running mirror pipeline Branch: cudnn-gh/pr-336-7ad1610 |
…ut names at lowering Two classic-parity items from the internal CI notebook set: - set_name after build is legal classic usage (sample 24 renames a tensor on an already-built graph): names are labels with no execution semantics, so _rename_tensor no longer consults the freeze — the label write bypasses the sealed-tensor guard explicitly, and the ambiguity policy still governs the name index. - User renames on op OUTPUTS now reach the lowered graph: push_output_attrs pushes the IR name, matching classic where the rename acts on the same object the cpp graph holds (visible in JSON dumps and wrapper.Graph canonical-name lookups). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Collaborator
Author
|
@cudnn-ci-bot run |
|
🚀 Running mirror pipeline Branch: cudnn-gh/pr-336-701d9fe |
…patched The internal tree layers a DSL engine by monkey-patching cudnn.pygraph lifecycle methods process-wide at import (cudnn.TBD). Under pytest-xdist any worker that collects those tests carries the patches into unrelated tests: signature introspection then sees the wrapper's (*args, **kwargs) and the patched create_execution_plans swallows the one-shot error (except Exception) — false negatives against pristine-class contracts. Detect the replacement via __qualname__ and skip LOUDLY with the reason, instead of failing on behavior that is not this class's. The proper fix remains scoping the internal patches (fixture install/uninstall) or excluding the TBD shard from the shared py_test run; these guards just make the contamination visible as skips rather than red. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Collaborator
Author
|
@cudnn-ci-bot run |
|
🚀 Running mirror pipeline Branch: cudnn-gh/pr-336-f8774d8 |
Merged
Anerudhan
added a commit
that referenced
this pull request
Aug 6, 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 NVIDIA/cudnn-frontend#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)
* remove unprofessional comments (#349)
* BSA: avoid guardword scanner false positives (#350)
* benchmark: fix repo-root path resolution in bench_moe (#348)
* Python-native cudnn.pygraph: graph IR + pluggable execution backends (#336)
* feat(python): backend-agnostic native graph + Router (unification proposal)
Modernize the Python-native graph API into the backend-dispatch architecture
from the Frontend v1 "Python API Engine and Graph API Unification" proposal.
Graph construction stays backend-agnostic; a backend is chosen by a first-class
Router at create_execution_plans() time (per Anerudhan's feedback), and the
backend-specific representation (e.g. the C++ cuDNN graph) is generated lazily
only then:
Python Graph API -> create_execution_plans() -> Router -> selected backend
(native engine, else cuDNN)
Layers kept separate:
- Graph IR (Node/Tensor/NativeGraph): engine-agnostic op DAG, full introspection
- BaseEngine: the backend contract (check_support/execute/get_workspace_size +
priority); cuDNN Graph is one routed backend, not a hardcoded default
- Router (engines/router.py): first-supporting by priority; None => cuDNN
Included: the IR, BaseEngine, Router, a CPU-only ReferenceMatmulEngine
(CI-testable correctness oracle), the optional MatmulCuTileEngine, and node
builders for block-scale / MoE / reduction so a DSL fusion backend can consume
them via graph.nodes (replacing the monkey-patch "recorder").
Deferred to follow-ups (see docs/python_native_graph_router.md):
NativeGraph.from_pygraph() (raises NotImplementedError for now), the DSL fusion
backend port, attention backends, and cuDNN lowering of the new node types.
Tests: 42 passing on CPU (IR + Router + reference-engine execute + cuDNN
fallback); cuTile path gated to SM100.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* refactor(python): trim NodeType to exercised ops; doc mixed candidate-list routing
- NodeType now lists only the op types this version exercises; drop the unused
norm/reshape/slice/etc. entries (re-add per-op when needed, following the
block-scale / MoE / reduction examples).
- Document the target routing model: create_execution_plans() takes one mixed
candidate list (native engines + cuDNN heur_modes) and produces a ranked list
of plans across backends; this PR ships the first-supporting-by-priority form.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* refactor(python): drop BATCHNORM / BATCHNORM_INFERENCE from NodeType (unused)
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* refactor(python): remove conv ops from native graph (unused foundation)
Drop CONV_FPROP / CONV_DGRAD / CONV_WGRAD: enum entries, the conv_fprop /
conv_dgrad builders, their dim inference in nodes.py, cuDNN lowering branches,
and the conv test. Re-add per-op when a backend needs conv, following the
block-scale / MoE / reduction examples.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs(python): use generic 'python DSLs' for backend examples
Avoid naming specific internal backends in public docs/docstrings.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* refactor(python): unify engines into one flat engine-id space (no cuDNN wrapper)
Replace the single-selected-backend + "if native else cpp" fork with the
engine-id model: python engines and cuDNN backend engines share one flat id
space. Python engines occupy a reserved high region (engine_ids.py,
PYTHON_ENGINE_ID_BASE = 1<<20) and each declares a stable engine_id it owns, so
ids never shift with registration order (reproducible autotune / pinned plans).
- engine_ids.py: PYTHON_ENGINE_ID_BASE + is_python_engine() + a phase-1
CUDNN_HEURISTIC_ENGINE_ID sentinel. Single source of truth for the namespace.
- Router.select()->one-engine becomes Router.plan()->ranked list of
PlanConfig(engine_id, knobs): supporting python engines (by id) + one trailing
cuDNN entry. TODO: interleave the true per-engine cuDNN configs
(get_engine_and_knobs_at_index) + real heuristics ranking; for now just concat.
- NativeGraph: _selected(engine) -> _plans(list) + _plan_index; add
get_execution_plan_count() / select_plan(i). check_support / build_plans /
get_workspace_size / execute all dispatch on the selected plan's id via
is_python_engine — one predicate, no fork. cuDNN is lowered lazily only when a
cuDNN-id plan is selected (pure-python when a python plan wins).
- BaseEngine: drop `priority`, add stable `engine_id` (reserved region).
reference_matmul = BASE+0, matmul_cutile = BASE+1.
Tests updated to assert the plan list; 41 pass on CPU incl. cuDNN fallback.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): make cudnn.pygraph engine-aware in place (transparent front door)
Users keep the classic API — g = cudnn.pygraph(...) is unchanged for every
existing sample — yet a graph transparently routes to a registered python engine
when it's fully represented. No new user-facing class, no rename.
pygraph_engines.install(pygraph) (called from __init__, same sanctioned pattern
as pygraph.execute = _execute) augments the pybind class in place:
- Per-graph mirror (WeakKeyDictionary) records a Node/Tensor IR alongside the
real C++ calls for a curated represented set (matmul + common pointwise),
mirrored via the NativeGraph builders so the recorded op is exactly what
engines consume.
- Every other op-builder is auto-wrapped to flag the graph "opaque" — the safe
direction: only disables the python path, never changes classic output.
- Lifecycle (create_execution_plans/check_support/build_plans/get_workspace_size/
execute/build) routes to a python engine iff one is registered AND the whole
graph is represented AND it supports the graph; else delegates to the untouched
C++ path.
Verified on an L40S against the real cuDNN build: a classic matmul runs
byte-identically with and without the augmentation, and a matmul+bias+relu graph
built via cudnn.pygraph + ReferenceMatmulEngine routes to the python engine with
exact results. Eager for now (C++ graph still built); lazy/pure-python is the
follow-up (needs a structured builder per op — multi-tensor returns like sdpa
can't be mirrored generically). NativeGraph stays as the standalone/greenfield
authoring object sharing the same IR + engines.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): native GEMM-family lowering + fix cuDNN execute path (phase 1)
Toward the native cudnn.pygraph migration (GEMM-family first). Make the native
build->lower->cuDNN execute path actually work end to end, and extend lowering
coverage to the GEMM family.
Fixes (all latent — the cuDNN execute path had never been GPU-tested):
- Thread the cuDNN handle: NativeGraph(handle=...) -> passed to the lowered
cudnn.pygraph so heuristics/build have a handle.
- Propagate the IR uid to the C++ tensor (was uid=-1 for auto tensors), so
execute()'s variant pack (keyed by IR uid) actually binds the buffers.
- POINTWISE lowering: the C++ pygraph has no generic pointwise(); dispatch on the
mode to the named ops (relu/gelu/sigmoid/tanh, add/mul/sub/div; add/mul also
cover bias/scale via broadcast).
Lowering coverage added: reduction, block_scale_dequantize, block_scale_quantize
(2 outputs), moe_grouped_matmul.
Validated on GPU (SM89): matmul and matmul+bias+relu built natively via
NativeGraph, lowered to cuDNN, execute with exact parity (new
test_native_cudnn_lowering.py, GPU-gated). Full native/router/pygraph suite: 45
passing. Per-op output-shape inference (e.g. reduction reduced dims) and
block-scale/moe execution parity are the next slices.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): reduction output-shape + SF reordering lowering (GEMM-family phase 2)
- reduction(): take an explicit reduced `dim` (cuDNN requires the reduction
output dims set); lowering sets set_dim/set_stride on the cuDNN op. Validated
matmul -> reduction(ADD over N) parity on GPU.
- lower_tensor(): propagate reordering_type to _make_tensor (e.g. F8_128x4),
needed for block-scale scale-factor tensors.
Native/router/pygraph + GPU parity suite: 46 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): native block-scale (nvfp4) lowering on Blackwell + fixes (phase 3)
Complete the GEMM-family native lowering with block-scale, validated on SM100.
Two more latent cuDNN-path bugs fixed:
- _lower_to_cpp passed io_data_type=None -> cudnn.pygraph rejects None. Now omit
io when unset; default intermediate/compute to FLOAT (matching cudnn.graph())
so cuDNN infers virtual (intermediate) tensor dtypes during build.
- lower_tensor now propagates reordering_type (F8_128x4) and omits data_type
when unset (NOT_SET) so cuDNN infers fused block-scale dequant output types.
Validated on SM100: dequant(A_fp4)@dequant(B_fp4) with F8_128x4 SFs builds +
executes via NativeGraph (test gated to SM100 + torch fp4; parity harness = the
repo's own fp4 test, which also only checks execution).
CPU overhead of the native Python layer (512^3 fp16, L40S): build +0.40 ms on
~106 ms (~0.4%, dominated by cuDNN heuristics); execute +0.3 us/call
(9.8 -> 10.1 us). Negligible.
Native/router/pygraph + GPU parity (matmul, bias+relu, reduction, block-scale):
48 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): native moe_grouped_matmul lowering + parity (GEMM-family complete)
- Add moe output-shape inference (token [1,T,H], weight [E,H,N] -> out [1,T,N])
so NativeGraph.validate() passes; cuDNN infers the same at build.
- GPU parity test (self-contained per-expert reference; no dependency on the
upstream test's helper) — validated on SM100.
GEMM family now fully native-lowered + validated on GPU: matmul, pointwise
(bias/relu), reduction, block-scale nvfp4, moe. Suite: 48 passing.
Next: non-GEMM ops (norms/reshape/slice/...) then the C++ _op rename + atomic
flip of cudnn.pygraph.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* fix(python): IR-uid -> C++-uid translation at execute; native rmsnorm (first norm)
Systemic fix: op-created C++ tensors (op outputs / virtuals) get uids assigned
by the C++ FE during build_operation_graph, in ITS enumeration order — which
does not match IR allocation order for multi-output ops (rmsnorm assigns
INV_VARIANCE=5, Y=6 while the IR allocated Y=5, inv_var=6). Keying the variant
pack by raw IR uids bound Y's buffer to inv_var: a [N,C,H,W] fp16 write into a
16-byte buffer (heap corruption / NaN). Single-output ops only worked by
allocation-order coincidence.
Fix: keep the lowering tensor_map; after build_operation_graph query every C++
tensor's real uid into an explicit IR-uid -> C++-uid map; execute() translates
variant-pack keys through it. No more order coincidence anywhere.
rmsnorm added as the first-class norm template (per "no corner-cutting" — the
generic opaque-op bridge was rejected/reverted since it makes non-GEMM ops
un-introspectable black boxes): named input/scale/epsilon/bias ports, Y/inv_var
outputs, norm_forward_phase param, pass-by-value epsilon; Y/inv_var dims carried
in the IR, cuDNN infers on its side. GPU parity: errY=0.0019, errI=0.0.
Suite: 49 passing (GEMM family re-validated through the translation path).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* refactor(python): Python IR owns the uid namespace end to end
Systematic uid review — four assignment paths existed:
1. user at creation: tensor(uid=...) (pybind _make_tensor, default -1)
2. user post-creation: tensor.set_uid() (mainline integrator pattern)
3. C++ FE auto-assign at build_operation_graph (enumeration order,
nondeterministic for multi-output ops) <- the coincidence trap
4. Python IR _alloc_uid (eager, sequential)
New invariant: for Python-built graphs, (3) NEVER triggers. The IR assigns
every uid eagerly at creation (auto or user-specified); lowering pushes ALL of
them explicitly to C++ — inputs via _make_tensor(uid=), op-created
outputs/virtuals via one set_uid loop over the complete tensor_map (single
point, impossible to forget per-op). Mixed construction (extending the lowered
C++ graph directly) is unsupported: a graph is pure-Python or pure-C++.
- Replace the IR->C++ uid translation map with a post-build ASSERTION: a
lowering path that fails to push a uid now fails loudly instead of being
silently translated (or worse, mis-binding buffers).
- _alloc_uid skips user-reserved uids; duplicate explicit uids rejected eagerly
at tensor() (C++ would only fail at build).
- execute() keys the variant pack by IR uids directly (== C++ uids by
construction).
Suite: 50 passing on SM100 (rmsnorm multi-output canary + block-scale included).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): full pointwise coverage — 54 ops, table-driven, mode == method name
Cover the entire pointwise surface of the C++ pygraph (54 methods) natively:
- Canonical op kind: params["mode"] IS the C++ pygraph method name (the
pointwise_mode enum is not exposed to Python; the method name is the semantic
name). Lowering collapses to a direct getattr dispatch — the mode<->method
mapping table is deleted as a concept.
- 47 uniform ops are generated from _POINTWISE_TENSOR_ARGS, a table of the
pybind tensor-argument names per op (mirrors the C++ signatures), so both
positional and the classic keyword call styles (bias(input=, bias=),
max(input0=, input1=)) work — required for the eventual cudnn.pygraph flip.
- 7 ops with scalar attributes get explicit builders storing them in params
(introspectable): relu(negative_slope/lower_clip/upper_clip), leaky_relu,
swish(swish_beta), gen_index(axis), + relu/leaky_relu/swish backwards.
Lowering forwards them as keywords.
- ReferenceMatmulEngine: keys move to method names; declines pointwise nodes
carrying scalar attributes it does not implement (correct-by-construction).
- Front-door mirror: classic calls passing scalar extras (e.g. relu clips) now
flag the graph opaque instead of silently dropping the attribute and
mis-routing to a python engine.
Tests: every builder exercised in both call styles + scalar-attr introspection
(CPU); sqrt/abs/max/min chain through real cuDNN on GPU. 53 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): norm family via one declarative table (10 ops, generic lowering)
All norms native — rmsnorm(_backward), layernorm(_backward), adalayernorm(_backward),
instancenorm(_backward), batchnorm, batchnorm_inference, batchnorm_backward —
through ONE mechanism instead of per-op code:
- _STRUCTURED_OPS: a declarative table per op — NodeType, tensor-input ports
(== the C++ pybind kwarg names), enum/scalar params (norm_forward_phase,
has_dbias), output ports in C++ return order, and per-output shape inference
(IR-side dims for introspection; cuDNN re-infers at build). Builders are
generated (keyword call style, as these ops are used repo-wide); lowering is
one generic branch: kwargs assembly + one call + zip outputs.
- List inputs (batchnorm peer_stats) become indexed ports (peer_stats_i) + a
count param, reassembled at lowering.
- The hand-written rmsnorm builder AND its lowering branch are deleted —
migrated into the table; the suite re-validates rmsnorm through the generic
path (multi-output uid canary intact).
GPU parity: layernorm fwd (Y/mean/inv_var) + layernorm_backward (DX/DScale/
DBias) vs torch autograd, using the supported LN config ([N,C,1,1]
channels_last, as in classic test_layernorm — the initial row-major 4D attempt
fails identically on the classic API, i.e. a kernel-support limit, not a
lowering bug). CPU: every table op builds a first-class node with named ports;
peer_stats port machinery covered. 56 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): conv + structural ops; collapse ALL structured ops into one table
_STRUCTURED_OPS now covers 25 ops — norms (11 incl. genstats), reduction,
block-scale (de)quantize, moe fwd/bwd, conv fprop/dgrad/wgrad, reshape, slice,
transpose, concatenate, rope fwd/bwd — one declarative entry each, one generic
lowering branch. Only matmul (positional ergonomics + front-door mirror) and
sdpa fwd/bwd (conditional kwarg assembly) remain explicit.
Deleted in the collapse: the hand-written reduction / block_scale_dequantize /
block_scale_quantize / moe_grouped_matmul builders AND their four lowering
branches, plus nodes.py moe shape inference (moved to the table). The suite
re-validates all of them through the generic path on GPU.
Table mechanics extended (each a one-word spec key, no new concepts):
- attrs: scalar/enum/list params forwarded verbatim (padding vectors, axis,
slices, permutation, reshape_mode, rope_dim, mode, ...). Conv accepts BOTH
the symmetric `padding` convenience and pre/post_padding — forwarded as
given; pybind overload resolution picks the right C++ binding.
- out_dims reserved kwarg (list, or {port: dims}): explicit output shapes for
ops cuDNN cannot infer — generalizes reduction's old `dim` param.
- push_output_dims: IR dims pushed to C++ for dgrad/wgrad/reduction/reshape/
moe_bwd (classic API also requires set_dim there).
- no_cdt: bindings without compute_data_type (reshape, concatenate).
- Builders accept tensors positionally or by port name; infer lambdas are
best-effort (try/except -> None; C++ validates at build).
GPU parity added: conv_fprop vs torch conv2d (NHWC), incl. asserting the
table's shape inference. CPU: all 25 ops x 2 call styles + out_dims. 58 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): sdpa family via generic kwarg capture — full ~130-arg surface
The six sdpa variants (sdpa, sdpa_backward, sdpa_fp8, sdpa_fp8_backward,
sdpa_mxfp8, sdpa_mxfp8_backward) are now declared in _CAPTURED_OPS, the third
and final table mechanism: builders capture ALL kwargs generically — tensor
values (incl. torch/dlpack) become named ports (port == C++ kwarg), scalars /
enums / score_mod callbacks go to params verbatim, dropout tuples are flattened
per element — and lowering rebuilds the kwargs for one C++ call. The full C++
kwarg surface (~130 args: paged attention tables, diagonal bands, sink tokens,
cu_seqlens, fp8 descales/amaxes, ...) is supported without hand-mirroring any
of it, and future binding args are picked up automatically.
Deleted: the explicit sdpa/sdpa_backward builders (~170 lines, common-args
only) + their two lowering branches + nodes.py sdpa shape inference (moved to
table lambdas — and fixed: O is q-shaped with v's head dim, not v-shaped).
Semantics now match the classic API exactly: sdpa always returns (O, Stats)
with Stats None in inference mode (generate_stats/is_inference logic); output
dim/stride are pushed to C++ (the SDPA node requires O's layout pre-validate —
that's how BSHD vs BHSD output is chosen).
GPU: sdpa causal fp16 EXECUTION parity vs torch SDPA (was build-only before).
59 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python)!: THE FLIP — cudnn.pygraph is now the Python graph class
The public cudnn.pygraph name now binds the Python IR class (class name:
pygraph; module: python/cudnn/pygraph.py — no "relative-to-history" naming).
The C++ graph builder is internal-only at cudnn._pybind_module.pygraph and is
reached exclusively through lowering: a graph is pure-Python or pure-C++,
never mixed. Zero C++ changes — the demotion is by namespace, not rebuild.
Deleted in the flip (afterthought residue):
- pygraph_engines.py front-door + its tests (no install()/monkey-patching
anywhere: register_backend is a native method on the class)
- NativeGraph.from_pygraph stub (meaningless now), use_native back-door
- docs/python_native_graph_router.md (initial-brainstorm doc, per review)
Drop-in surface for classic parity, driven by iterating the repo's own test
files until green (each item below was a real failure caught and fixed):
- conditional outputs ("maybe"): rmsnorm_backward(has_dbias=False) -> DBias
None; norm fwd INFERENCE -> mean/inv_var None; batchnorm next_running_*
present iff in_running_* given (classic returns None for absent outputs)
- torch interop: tensor(dim=x.size()) (torch.Size), data_type=torch.bfloat16
(converted at the C++ boundary via _library_type, IR stores user's value)
- output dtype semantics: an output without explicit set_data_type gets io
dtype (was mis-defaulted to intermediate FLOAT -> fp32 into fp16 buffers)
- Tensor gains the classic setter/getter surface (set_ragged_offset,
set_reordering_type, set_is_pass_by_value, ...); tensor_like(cudnn tensor);
tensor_scalar; CPU tensor_like -> pass-by-value (classic rule)
- ragged (THD) output layout: outputs' ragged_offset now pushed to C++ at all
mapping sites (was silently dense -> wrong values in sdpa_thd)
- validate-time table shape inference (topological): chained ops whose inputs
are virtual (conv on a relu output) infer once inputs are known;
builder-time infer stays as best-effort for direct inputs
- classic lifecycle: build_operation_graph lowers eagerly when no python
engines are registered, so deselect_*/query methods work between classic
steps via __getattr__ delegation to the lowered graph; build_plans(policy)
passthrough; deserialize(*args, **kwargs) passthrough incl.
enforce_precompiled; execute override_uids/shapes/strides + dlpack pointers;
get_execution_plan_count = python engines + backend's dynamically-queried
count (frontend NEVER statically enumerates backend engines — they vary by
backend version; Router keeps ONE delegating cuDNN entry by design)
- stride optional after set_dim (row-major inferred), None variant-pack keys
tolerated, C++-tensor keys resolved via get_uid
Validated: our suite (56) + classic spot-runs all green on real GPUs —
matmul_bias_relu, rmsnorm, layernorm, batchnorm, conv_fprop (incl.
execute_plan_at_index), apply_rope, kernel_cache, sdpa_with_caching, sdpa_thd,
sdpa_chunked_prefill (ragged+paged), conv_genstats, conv_reduction, slice,
block_scale_quantize_dynamic_shape, wgrads. Full-suite runs on SM100 + mhas in
flight; residuals to follow. Known pre-existing env skew (fails identically on
the unflipped installed package): test_deviceless_aot_compilation on this box.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): classic validate() timing + omit unset compute_data_type
Two classic-parity fixes surfaced by the full mhas run (3567 uniform failures,
one root cause):
- cudnnGraphNotSupportedError must fire at graph.validate(): the classic test
waiver pattern is try/except-skip AROUND validate(), with
build_operation_graph() called bare. With no python engines registered,
validate() now lowers and runs the C++ validate right there (unsupported
configs skip, not fail); build_operation_graph()/plan creation are staged
behind flags so each C++ step runs exactly once in classic sequencing.
Python-engine graphs still never touch C++ at validate.
- compute_data_type=None is now OMITTED at every lowering site (matmul /
pointwise / structured / captured) instead of passed through: classic ops
default to NOT_SET in C++; pybind rejects None. Also converts via
_library_type when set (torch dtype parity).
Previously-failing mhas case now skips as on classic; our suite 56 passing.
Full-suite + full-mhas reruns in flight.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* docs(router): codify the extension contract for the future heuristics MR
Ranking policy is intentionally undecided; what IS decided: policy pluggable at
three levels (Router subclass / per-graph / process default); plan() may return
any ordering or mix; backend engine sets are discovered per graph at plan time
(never statically enumerated); PlanConfig can carry concrete backend engine
configs, with pygraph._lower_cudnn_plan as the designated point to honor them.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): plan-selection lifecycle + registration validation (review items 2, 6)
Review item 2 (reproduced bugs):
- ONE plan index space: [0, n_python) are python plans, [n_python, ...) are the
backend's plans (sub-index = index - n_python, queried dynamically).
get_execution_plan_count() and select_plan() now agree; selecting a backend
sub-index lowers on demand, bui…
YangXu1990uiuc
added a commit
to YangXu1990uiuc/cudnn-frontend
that referenced
this pull request
Aug 7, 2026
docs/python_graph_and_execution_backends.md described the shape this PR replaces, so a reader following it would have looked for anchors and closed_under that no longer exist. Reviewer asked for a design doc; this one has been checked in since NVIDIA#336 and just needed to catch up. Four sections added -- what the manifest decides and what it deliberately leaves to the engine, how facts are attached and why after the freeze, how handle/stream/device actually resolve, and which imports each dispatch stage may pay for. Existing sections corrected where the PR moved them: ImportError joins the decline types, engines no longer declare their own id, register_backend is out-of-tree only. The stream question specifically: _resolve_stream is a fallback for having NO handle, not for a handle whose stream could not be read -- that raises, since running on the wrong stream is a correctness bug rather than a degradation.
3 tasks
YangXu1990uiuc
added a commit
that referenced
this pull request
Aug 7, 2026
* Scope engine dispatch to families, and hang facts off the graph
Four changes, all narrowing an over-broad promise:
1. _GEMM_CLOSURE no longer lists RESHAPE. closed_under is a PROMISE that the
family serves every node type in it, and nothing in gemm/frost consumes a
RESHAPE node. With it there, `matmul -> reshape` matched, only the matmul
was compiled, and execute failed with "the variant pack is missing buffers
for ['mm::C']" -- with no backend fallback, because the python engine had
already claimed the graph. It now declines at build with the honest "no
engine proposed a plan". Plain matmul routing is unchanged.
2. cu_seq_len_q / cu_seq_len_kv become a FACT (has_cu_seq_len) judged by a
Capabilities row, instead of poisoning SdpaGraphFacts.invalid. `invalid`
means malformed-for-everyone; putting a not-implemented-here feature there
would also bar the engine that eventually implements prefix sums. Net
eligibility is unchanged: no row sets cu_seq_len=True.
3. EngineRow -> EngineFamily. A family is the unit everything is scoped to:
node-type envelope, id block, arch range, maturity gate, facts vocabulary.
Splitting one (fp8 SDPA out of SDPA) costs one more entry and nothing
elsewhere. Adds the id-block disjointness check that was assumed but never
tested -- a shared id would make the engine an autotune result names
ambiguous.
4. Facts hang off the graph, scoped per family, and no caller asks for them.
validate() walks the manifest, finds the families that claim this graph,
runs each one's declared analyzer once and attaches the record; a graph no
family claims carries no payload. Engines read that record back instead of
parsing again. Facts are family-scoped by construction -- an SDPA fact
means nothing to a GEMM engine -- so this is a mapping, never a union
record every family would have to widen.
The record is keyed by the analyzer callable itself, so the ranking (which
resolves it from EngineFamily.analyzer before any engine module is
imported) and the engine (which passes the callable it already imports)
reach ONE record with no family-name string to keep in sync. That is the
drift the backend's SDPA heuristics have, where the feature vector and the
engine's own view of the graph are extracted separately.
Both SDPA families go through it. The bwd family was reading the same
module-private weakref cache as fwd, so this is not removing a duplicate
parse -- it is keeping the sharing after the cache moved onto the graph,
and making it reachable by the ranking rather than only by engines. That
weakref cache is gone and analyze() is now pure.
Attaching at validate() exposed a live staleness hole: build_operation_graph
-> _sync_ir_shapes_from_backend rewrites IR dim/stride with the backend's
inference (channels-last conv and friends), and facts describe layout. The
node count is unchanged, so nothing would have caught it; the records are
now dropped and re-attached there. Lazy evaluation had hidden this by
always running after the sync.
Verified on Blackwell (sm100): test_engine_router 62, sdpa graph-analyzer 66.
* Make a family a kind of graph, and stop paying for the DSL to decline
The manifest had two different things wearing one name: three entries owned a
100-wide id block and returned several engines, five owned a single id and
returned one. GDN and KDA graphs matched two entries each, so "which family
serves this graph" had no single answer.
A family is now a KIND OF GRAPH -- roughly what the backend calls an
operation-graph mode, at a granularity of our choosing -- and every graph
belongs to exactly one or to none:
gdn {GDN, GDN_BWD} GdnFrostEngine + GdnCuTileEngine one block
kda {KDA, KDA_BWD} KdaFrostEngine + KdaCuTileEngine one block
gdn2, frost_gemm, frost_sdpa_fwd, frost_sdpa_bwd
Classification is a lookup (_FAMILY_OF_NODE), not N families each declaring a
claim that then has to be proven disjoint: a function returns one value, so
"two families claimed this graph" is not a case that can arise and is not an
invariant anyone has to test. A graph naming two families (a matmul and an sdpa
together) belongs to neither and goes to the backend.
Each family reserves FAMILY_BLOCK ids and never opens a second block, so
engine_id alone identifies the family. Ids are pre-release (engine_ids.py), so
the blocks are re-cut here; the earlier claim that fp8 SDPA could not be split
out because its ids had shipped was simply wrong.
closed_under is deleted. It existed to reject a graph without paying for an
import, and it duplicated a judgment the engine has to make anyway -- which is
how it came to promise RESHAPE support that nothing implemented. The real
defect was in the gemm analyzer: _node_to_recorded_op returned None for an
unrecognized node and the caller SKIPPED it, so any unhandled node type
silently compiled a subgraph and execute then demanded buffers the caller never
bound. It declines now.
That deletion is only safe because declining no longer costs an import. Support
checks were dragging the CuTe DSL in: cudnn/sdpa/__init__ eagerly imported .bwd
and .fwd, and fwd/engines.py -- which holds Capabilities and mismatch, both pure
data -- imported api_dsl at module level to bind EngineSpec.lower. The three
package inits are lazy (PEP 562, the pattern cudnn/__init__ already used) and
the adapter resolves at build time:
import cudnn.sdpa.graph_analyzer 1059 ms, +381 modules -> 9.4 ms, +2
support-check module +1387 modules -> +7
Facts and capabilities now speak cudnn.data_type instead of torch.dtype. Facts
are what every engine of a family reads, so expressing them in one framework's
types would make dispatch require that framework; torch appears only where a
torch tensor is actually allocated (graph_analyzer.to_torch_dtype).
Finally, planning does finalize -> freeze -> analyze in that order, and the
finalize runs the SAME way whether or not an out-of-tree engine was registered.
Lowering to C++ and reflecting layout back is how the graph learns the strides
it will execute with -- a property of the graph, not of whichever engine serves
it. Splitting those paths is what made the frozen snapshot unenforceable: the
registered-engine path lowered later, inside the Router, and
_sync_ir_shapes_from_backend writes through object.__setattr__ specifically to
bypass the freeze. _facts_for() memoizes only a frozen graph, so there is no
invalidation rule left to get wrong.
heuristics_sort still does not read facts. That is deliberate and now says so:
the seam is in place so that writing a real policy -- order a family's engines
on its facts, then merge against the backend on predicted time -- does not mean
re-plumbing the graph first.
Verified on Blackwell sm100: test_engine_router 64, sdpa graph-analyzer 66,
sdpa/frost + gemm/frost 4500 passed / 2119 skipped at -n 32, identical to the
counts before this change. linear_attention fails 286 tests here and on
pristine gh/develop alike (CUDA graph capture), untouched by this.
* Take the device probe off torch, and fix what the third review found
Facts no longer read torch.cuda. Compute capability and SM count come from
cudnn.create_device_properties() -- the backend's OWN device descriptor, the
same object the C++ deviceless-AoT path serializes and replays. graph_analyzer
now imports no torch at module level at all; torch appears only where a torch
tensor is actually allocated. This is also the step that makes a deviceless
python engine possible: facts computed from a serialized descriptor need no
live device, whereas torch.cuda.current_device() required one.
frost/buffers.py gains current_device_id() with the same two-probe shape as
current_sm(): a missing cuda-python must not look like a missing GPU.
Review fixes:
- _ATTACHABLE was defined, unused, and its comment described behaviour that
did not exist. Deleted.
- The three lazy __init__ files returned only __all__ from __dir__(), hiding
every normal module attribute including __name__. Union with globals() now.
- linear_attention/frost/__init__ eagerly imported the GDN, KDA and GDN2
engines, so importing one family's engine pulled its neighbours -- which
defeated the per-engine ImportError tolerance the family factories exist to
provide. Lazy now.
- _facts_for() keyed on f"{module}.{qualname}", contradicting its own
docstring and letting a reloaded same-named callable collect another
analyzer's record. Keyed on the callable itself.
- frost/README.md still described backend-first ranking, including the
pseudocode; the implementation has been python-first since the seam landed.
Verified on Blackwell sm100: test_engine_router 64, sdpa graph-analyzer 66,
sdpa/frost + gemm/frost 4500 passed / 2119 skipped at -n 32 -- the same counts
as before this series began.
* Make the manifest the single source of engine ids, and test the import boundaries
Engine ids were assigned in four places: two _ID_OFFSETS tables (sdpa fwd and
bwd) and class attributes on the gemm and linear-attention engines. The manifest
only VALIDATED containment after the fact, so "two engines share a slot" was
possible in a way "two families overlap blocks" was not.
The manifest now assigns. Each family lists its engines as slots:
slots={"sdpa_fwd_prefill_sm100_d128": EngineSlot(0, opt_in=True), ...}
and instantiate() hands the factory {name: engine_id}. Engines carry no id of
their own, so one cannot claim a number it was not given -- the error stops
being caught and starts being unrepresentable. The whole python id space reads
out of one file instead of being reconstructed from four.
opt_in moves with it, from per-family to per-ENGINE. Maturity is a property of
one implementation: the half-precision SDPA engines can now graduate while the
fp8 engine is still maturing, which one flag per family made impossible. It
stays in the manifest rather than on the engine class because the whole point
of the gate is to know what to offer WITHOUT importing the engine.
register_backend() is now only what its docstring always claimed. In-tree
engines never registered -- the manifest discovers them -- but eleven tests
still called it, and an "the in-tree owner may register itself" exemption
existed to keep that working. Every one of those calls was redundant (the
cuTile-declines test already asserts against the ranked plan list by name), so
they are gone, and with them the exemption, its test, and the namespace-
containment patch this series had added to keep it alive. The check is now one
rule: ids below OUT_OF_TREE_ID_BASE are rejected.
Two tests replace what was runtime luck:
- test_every_engine_spec_has_a_manifest_slot: an engine added without a slot
would silently never be built. Checked both ways, plus slot uniqueness and
range, on CPU.
- test_import_boundaries.py: what each dispatch stage may drag in. The graph
API must not require a framework, and deciding whether an engine COULD serve
a graph must not import the machinery that would serve it -- deleting
closed_under is only safe while that holds. Each check runs in a FRESH
interpreter (a module imported by the test process would make an in-process
assertion pass for the wrong reason) and measures the DELTA against an empty
one (nvidia_cutlass_dsl is injected at startup by a .pth, so an absolute
check blames us for what the interpreter did first).
Making that pass took one more step: fwd/engines.py and bwd/engines.py still
imported torch at module level for lowering helpers that share the file with
Capabilities and mismatch(). Deferred, so the support-check modules now pull
110 modules instead of 1103, and neither torch nor cutlass.
Verified on Blackwell sm100: test_engine_router 64 + test_import_boundaries 5,
sdpa graph-analyzer 66, sdpa/frost + gemm/frost 4568 passed / 2119 skipped
at -n 32.
* Fix the two review findings that still apply
to_torch_dtype() indexed its map directly, so an unmapped cudnn.data_type
raised KeyError deep in lowering. Only Q's dtype is capability-checked, and
tensor_desc_from_ir() runs on every bound tensor, so O / Stats / a side output
can still carry one. It declines now.
Two docstrings still named validate() as the site that runs the analyzer;
planning does, after the freeze.
* Classify a graph without asking which machine is asking
family_for() took an `sm` and folded availability into its answer, so None
meant either "not that kind of graph" or "no engine for it here" and a caller
could not tell which. What kind of graph something is cannot depend on the
machine: classification is now pure, and EngineFamily.offered_ids(sm) answers
availability separately. A matmul graph is a gemm graph on a host with no gemm
engine at all.
_FAMILY_OF_NODE becomes _ANCHOR_NODE_TO_FAMILY. The old name read as "every
node maps to a family", when the table holds only the node types that NAME
one -- POINTWISE, REDUCTION, a type added tomorrow are absent on purpose and
ignored, which is why `matmul + pointwise` is a gemm graph. Whether a family
can serve the WHOLE graph stays its analyzer's judgment; a coarser copy of
that here is what closed_under was. The comment now says so at the table
rather than leaving it to be discovered.
test_classification_does_not_depend_on_the_machine pins it: flipping the
opt-in flag or asking about another arch changes what is OFFERED and not what
the graph IS.
Verified on Blackwell sm100: test_engine_router 65 + test_import_boundaries 5,
sdpa/frost + gemm/frost 4500 passed / 2119 skipped at -n 32.
* Declare an architecture line, not the parts of it that exist today
Capabilities.arches was an exact set -- frozenset({(10, 0), (10, 3)}) -- checked
with `facts.device_cc not in capabilities.arches`. An sm100 kernel runs on the
whole sm100 line, so enumerating the members that exist today silently declines
the ones that ship later: Rubin (sm107) and Thor (sm110) are meant to reuse
these kernels and both were excluded. It is a range now, inclusive, encoded
major*10 + minor as engines/manifest.py already did, and the decline message
says "requires SM100-119" instead of listing device families.
The manifest's own sm_lo/sm_hi is deleted, for the reason closed_under was:
a coarser duplicate of a judgment the engine has to make anyway, which is a
second thing to maintain and a place to lie. It already lied, and in exactly
the way that matters here -- frost_gemm capped at SM103 while kernel_registry
declares PIPELINE_ARCH_RANGES["sm100"] = ((100, 120),) with the comment "sm100
templates use only family-portable Blackwell instructions"; frost_sdpa_bwd
capped at SM121 against a Capabilities row of 120-129. On Rubin the gemm family
would not have been offered at all, so its engine never got asked.
Deciding an engine is wrong for a device is now said once, by the engine.
Dropping the manifest copy costs one module import before a decline, which the
laziness work already made cheap: measured on an sm90 host, the SDPA family
instantiates its seven engines and they decline, pulling no CuTe DSL
(test_import_boundaries.py holds that). current_sm() leaves the dispatch path
entirely -- engines_for(graph) and _attach_facts no longer probe the device,
so there is one less way for device state to reach classification.
The SDPA test suites had five hand-copied _is_sm100 gates, every one pinned to
exactly (10, 0), so all five skipped on sm103 while the engines they test serve
the line; only the MXFP8 file had it right. conftest.py now carries one
requires_blackwell / requires_blackwell_geforce / requires_dsl, aligned with
what the engines declare rather than re-derived per file.
Verified on Blackwell sm100: test_engine_router 65 + test_import_boundaries 5,
sdpa/frost + gemm/frost 4500 passed / 2119 skipped at -n 32.
* Unify the SDPA suite's run gates, and trim the comments this series added
Eight test files each defined their own arch and DSL gates -- five copies of
_is_sm100 all pinned to exactly (10, 0), and _dsl_deps_available /
_dsl_available / _require_dsl for one `import cutlass`. They now come from
frost_test_utils, matching the convention gemm_test_utils already set, so the
gate is stated once and against what the engines declare.
Comments trimmed to the load-bearing facts; the measurements and the history
behind them are in the commit messages of this series where they belong. One
of them had already gone stale -- EngineFamily's docstring still listed "the
arch range" among what a family owns, one commit after that was deleted.
* Decline a missing or too-old CuTe DSL at check time, not at codegen
Deferring the DSL import past check_support moved a failure from discovery to
execution, and the two stages do not catch the same things. Without the cutedsl
extra installed, check_support passed (it reads Capabilities and facts only),
then build_plan raised ImportError -- which is in neither the engine's except
list nor decline_types(), so build_plans() propagated it instead of walking on.
The backend was in the same ranked list and never got its turn. Before this
series the DSL was imported at module scope, instantiate() caught the
ImportError, and the family simply vanished; the laziness work broke that
without replacing it.
Two answers, in the right order:
- The engine now DECLINES at check_support, using probes that do not execute
the module: importlib.util.find_spec("cutlass") (7 ms, 30 modules) and
importlib.metadata.version (5 ms, 2), against ~4.3 s and 1410 modules for the
real import. A plan that cannot be lowered no longer enters the ranked list
at all.
- ImportError joins decline_types() and the engines' build-time except clause,
as the second line rather than the only one.
The CuTe primitives these engines lower through need 4.7.0, so the version is
part of the same check -- an older DSL fails during codegen naming a missing
attribute rather than a version. Two wrinkles that cost a first attempt:
pyproject declares nvidia-cutlass-dsl while this box had
nvidia-cutlass-dsl-internal, so a single-name lookup finds nothing on one of
the two; and internal RCs number themselves independently ("0.3.0+2026..."),
so the public floor cannot judge them -- my first version declined the very
build every test here had been passing on. Unparsable, absent, and internal all
count as not-too-old: refusing on a string we failed to read would reject a
machine that works.
Verified on the PUBLIC wheel, which nothing had run against before: swapped
nvidia-cutlass-dsl-internal 0.3.0 for nvidia-cutlass-dsl 4.7.0 and got
4500 passed / 2119 skipped on sm100, identical to the internal RC.
* Declare the CuTe DSL floor the engines actually require
The cutedsl extra asked for >=4.5.0 while the primitives these engines lower
through need 4.7.0, so pip would install a version support checks then decline.
CUTEDSL_MIN_VERSION and this bound are the same number now.
* Bind torch before the branch that needs it
Deferring torch put `import torch` inside one branch of the execute path while a
later, independent branch also used it: an fp8 graph takes the second without
the first, so `carver.take(facts.b, torch.int32)` raised UnboundLocalError. 30
tests in test_mhas_v2 failed on it.
Only that suite could see it. sdpa/frost pins an engine and exercises the kernel
directly; test_mhas_v2 goes through routing, which is what reaches
synth_kv_padding. The other three deferred imports are unconditional at function
scope and were checked.
test_mhas_v2 on sm100: 2178 passed / 707 skipped, and FROST serves 413/3201
graphs across four engines (d128 182, d256 122, d128_fp8 55, d192_d128 54) --
so family consolidation, manifest-assigned ids, the arch range and the DSL
laziness leave routing intact.
* Keep the DSL floor out of the dependency, where it would be contagious
Pinning the cutedsl extra to >=4.7.0 (the previous commit) would make
cudnn-frontend incompatible with anything holding the DSL back -- quack-kernels
pins ==4.6.0, and vLLM and friends carry their own constraints. Reverted to
>=4.5.0, with the reason at the pin rather than in a commit nobody will find.
Reviewer flagged the same thing.
The 4.7.0 floor is real, so it lives at runtime where it costs only the engines
that need it: check_support() declines an older DSL and the graph goes to the
backend. The test gate now uses the same predicate and the same constant, so an
older DSL makes the SDPA suite SKIP rather than fail -- it had checked presence
only, which would have run the tests and let them fail for a reason the suite
already knew.
sm100: sdpa/frost + gemm/frost + test_mhas_v2 = 6678 passed / 2826 skipped,
FROST serving 413/3201 graphs.
* Bring the design doc up to what dispatch now does
docs/python_graph_and_execution_backends.md described the shape this PR
replaces, so a reader following it would have looked for anchors and
closed_under that no longer exist. Reviewer asked for a design doc; this one
has been checked in since #336 and just needed to catch up.
Four sections added -- what the manifest decides and what it deliberately
leaves to the engine, how facts are attached and why after the freeze, how
handle/stream/device actually resolve, and which imports each dispatch stage
may pay for. Existing sections corrected where the PR moved them: ImportError
joins the decline types, engines no longer declare their own id,
register_backend is out-of-tree only.
The stream question specifically: _resolve_stream is a fallback for having NO
handle, not for a handle whose stream could not be read -- that raises, since
running on the wrong stream is a correctness bug rather than a degradation.
YangXu1990uiuc
added a commit
to YangXu1990uiuc/cudnn-frontend
that referenced
this pull request
Aug 11, 2026
graph.execute() inspected the caller's buffers twice and differently. The
backend path built a {uid: pointer} dict (_native_var_pack), whose _ptr
accepted a bare device address. A python engine got the caller's objects
untouched via resolve_node_buffers and reached them through
frost.buffers.probe, which raised "buffer of type int exposes neither
__cuda_array_interface__ nor __dlpack__" for that same address. One public
call, two answers, and the caller does not choose which plan the heuristics
land on.
Normalization now happens once, at the top of execute(), into Operands: the
caller-filled uids ascending, a ctypes pointer array, and — when a python
engine will read them — a Tensor record per operand carrying the buffer's own
dim/stride/data_type. Below that line the backend takes ctypes.addressof(ptrs)
and every engine takes pointers plus records. A bare address that the backend
took now reaches an engine too, shaped by the geometry the graph declares for
that port.
The order comes from exactly one source, never a union: the lowered graph's
variant-pack template when there is one (only C++ can see every user slot — a
tensor's ragged_offset is an operand but hangs off the Tensor rather than off a
node port, and the slots the graph fills itself must be excluded), and the IR
only for the python-only ops that cannot lower at all. The two sides never have
to agree: each indexes the layout it was handed.
C++ already turned a uid map into sorted pointers internally
("uid map -> extract sorted ptrs, delegate to the sorted_ptrs implementation",
graph_interface.h), so passing the array directly drops one dict build here,
one map copy in pybind and one hash lookup per operand there.
execute_with_raw_ptrs gains a plan_index because it only ever ran
plans.candidate, which stops being the plan the python walk built once the walk
has skipped an entry; the vector overload it duplicated had no callers and goes.
The pointer array is allocated PER CALL. Two threads may execute one graph
concurrently with different buffers, and a shared array hands each thread the
other's pointers — silently, since each pointer in it is individually valid.
The new test fails with [0,2,7,0,1,2,2,14] crossed results when the array is
shared.
Also deleted: the 87-line execute/execute_plan_at_index pair monkey-patched
onto backend_graph in __init__.py, unreachable since NVIDIA#336 made cudnn.pygraph a
python class that defines both names itself; the two always-false
`hasattr(graph, "_execute_with_ptrs")` fast paths in experimental/ops/sdpa.py
and the uid_order cache feeding them; and the five places
docs/adding_torch_custom_ops.md told authors to hand-roll that path, which
raises AttributeError as written.
Backend execute on a 128^3 bf16 matmul: 16.17 -> 14.76 us.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
YangXu1990uiuc
added a commit
to YangXu1990uiuc/cudnn-frontend
that referenced
this pull request
Aug 11, 2026
graph.execute() inspected the caller's buffers twice and differently. The
backend path built a {uid: pointer} dict (_native_var_pack), whose _ptr
accepted a bare device address. A python engine got the caller's objects
untouched via resolve_node_buffers and reached them through
frost.buffers.probe, which raised "buffer of type int exposes neither
__cuda_array_interface__ nor __dlpack__" for that same address. One public
call, two answers, and the caller does not choose which plan the heuristics
land on.
Normalization now happens once, at the top of execute(), into Operands: the
caller-filled uids ascending, a ctypes pointer array, and — when a python
engine will read them — a Tensor record per operand carrying the buffer's own
dim/stride/data_type. Below that line the backend takes ctypes.addressof(ptrs)
and every engine takes pointers plus records. A bare address that the backend
took now reaches an engine too, shaped by the geometry the graph declares for
that port.
The order comes from exactly one source, never a union: the lowered graph's
variant-pack template when there is one (only C++ can see every user slot — a
tensor's ragged_offset is an operand but hangs off the Tensor rather than off a
node port, and the slots the graph fills itself must be excluded), and the IR
only for the python-only ops that cannot lower at all. The two sides never have
to agree: each indexes the layout it was handed.
C++ already turned a uid map into sorted pointers internally
("uid map -> extract sorted ptrs, delegate to the sorted_ptrs implementation",
graph_interface.h), so passing the array directly drops one dict build here,
one map copy in pybind and one hash lookup per operand there.
execute_with_raw_ptrs gains a plan_index because it only ever ran
plans.candidate, which stops being the plan the python walk built once the walk
has skipped an entry; the vector overload it duplicated had no callers and goes.
The pointer array is allocated PER CALL. Two threads may execute one graph
concurrently with different buffers, and a shared array hands each thread the
other's pointers — silently, since each pointer in it is individually valid.
The new test fails with [0,2,7,0,1,2,2,14] crossed results when the array is
shared.
Also deleted: the 87-line execute/execute_plan_at_index pair monkey-patched
onto backend_graph in __init__.py, unreachable since NVIDIA#336 made cudnn.pygraph a
python class that defines both names itself; the two always-false
`hasattr(graph, "_execute_with_ptrs")` fast paths in experimental/ops/sdpa.py
and the uid_order cache feeding them; and the five places
docs/adding_torch_custom_ops.md told authors to hand-roll that path, which
raises AttributeError as written.
Backend execute on a 128^3 bf16 matmul: 16.17 -> 14.76 us.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
YangXu1990uiuc
added a commit
to YangXu1990uiuc/cudnn-frontend
that referenced
this pull request
Aug 11, 2026
graph.execute() inspected the caller's buffers twice and differently. The
backend path built a {uid: pointer} dict (_native_var_pack), whose _ptr
accepted a bare device address. A python engine got the caller's objects
untouched via resolve_node_buffers and reached them through
frost.buffers.probe, which raised "buffer of type int exposes neither
__cuda_array_interface__ nor __dlpack__" for that same address. One public
call, two answers, and the caller does not choose which plan the heuristics
land on.
Normalization now happens once, at the top of execute(), into Operands: the
caller-filled uids ascending, a ctypes pointer array, and — when a python
engine will read them — a Tensor record per operand carrying the buffer's own
dim/stride/data_type. Below that line the backend takes ctypes.addressof(ptrs)
and every engine takes pointers plus records. A bare address that the backend
took now reaches an engine too, shaped by the geometry the graph declares for
that port.
The order comes from exactly one source, never a union: the lowered graph's
variant-pack template when there is one (only C++ can see every user slot — a
tensor's ragged_offset is an operand but hangs off the Tensor rather than off a
node port, and the slots the graph fills itself must be excluded), and the IR
only for the python-only ops that cannot lower at all. The two sides never have
to agree: each indexes the layout it was handed.
C++ already turned a uid map into sorted pointers internally
("uid map -> extract sorted ptrs, delegate to the sorted_ptrs implementation",
graph_interface.h), so passing the array directly drops one dict build here,
one map copy in pybind and one hash lookup per operand there.
execute_with_raw_ptrs gains a plan_index because it only ever ran
plans.candidate, which stops being the plan the python walk built once the walk
has skipped an entry; the vector overload it duplicated had no callers and goes.
The pointer array is allocated PER CALL. Two threads may execute one graph
concurrently with different buffers, and a shared array hands each thread the
other's pointers — silently, since each pointer in it is individually valid.
The new test fails with [0,2,7,0,1,2,2,14] crossed results when the array is
shared.
Also deleted: the 87-line execute/execute_plan_at_index pair monkey-patched
onto backend_graph in __init__.py, unreachable since NVIDIA#336 made cudnn.pygraph a
python class that defines both names itself; the two always-false
`hasattr(graph, "_execute_with_ptrs")` fast paths in experimental/ops/sdpa.py
and the uid_order cache feeding them; and the five places
docs/adding_torch_custom_ops.md told authors to hand-roll that path, which
raises AttributeError as written.
Backend execute on a 128^3 bf16 matmul: 16.17 -> 14.76 us.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
YangXu1990uiuc
added a commit
that referenced
this pull request
Aug 12, 2026
…k producer (#547) * Keep frozen-ness on the graph, where it is one flag _freeze() stored a _frozen flag on the graph AND on every Tensor, every Node and the GraphContext, and gave the latter three a __setattr__ guard so a direct attribute write would raise. Freezing is a property of the graph; four copies of the state, and three guards to read them, is not what enforcing it needs. The guards were also expensive in a way nothing measured. A dataclass __init__ assigns field by field, so overriding __setattr__ turns construction into one python-level call plus one failed `getattr(self, "_frozen", False)` lookup PER FIELD. Tensor has fifteen. Measured: 2.53 us to construct a Tensor, of which 2.14 us was the guard, on an object that is by definition not yet frozen. A graph pays it once per tensor and once per node, every build. What actually closes the mutation routes is unchanged: _check_mutable guards every setter and op builder, node.inputs/outputs/params become MappingProxy views, and dim/stride become tuples. Those are structural — they cost nothing per call and they cannot be bypassed. What is no longer an error is assigning `t.dim = [...]` directly on a frozen graph, which was never a route the API offered; the test now pins the routes it does offer. Tensor construction: 2.53 -> 0.56 us. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * Normalize the variant pack once; both paths read the same operands graph.execute() inspected the caller's buffers twice and differently. The backend path built a {uid: pointer} dict (_native_var_pack), whose _ptr accepted a bare device address. A python engine got the caller's objects untouched via resolve_node_buffers and reached them through frost.buffers.probe, which raised "buffer of type int exposes neither __cuda_array_interface__ nor __dlpack__" for that same address. One public call, two answers, and the caller does not choose which plan the heuristics land on. Normalization now happens once, at the top of execute(), into Operands: the caller-filled uids ascending, a ctypes pointer array, and — when a python engine will read them — a Tensor record per operand carrying the buffer's own dim/stride/data_type. Below that line the backend takes ctypes.addressof(ptrs) and every engine takes pointers plus records. A bare address that the backend took now reaches an engine too, shaped by the geometry the graph declares for that port. The order comes from exactly one source, never a union: the lowered graph's variant-pack template when there is one (only C++ can see every user slot — a tensor's ragged_offset is an operand but hangs off the Tensor rather than off a node port, and the slots the graph fills itself must be excluded), and the IR only for the python-only ops that cannot lower at all. The two sides never have to agree: each indexes the layout it was handed. C++ already turned a uid map into sorted pointers internally ("uid map -> extract sorted ptrs, delegate to the sorted_ptrs implementation", graph_interface.h), so passing the array directly drops one dict build here, one map copy in pybind and one hash lookup per operand there. execute_with_raw_ptrs gains a plan_index because it only ever ran plans.candidate, which stops being the plan the python walk built once the walk has skipped an entry; the vector overload it duplicated had no callers and goes. The pointer array is allocated PER CALL. Two threads may execute one graph concurrently with different buffers, and a shared array hands each thread the other's pointers — silently, since each pointer in it is individually valid. The new test fails with [0,2,7,0,1,2,2,14] crossed results when the array is shared. Also deleted: the 87-line execute/execute_plan_at_index pair monkey-patched onto backend_graph in __init__.py, unreachable since #336 made cudnn.pygraph a python class that defines both names itself; the two always-false `hasattr(graph, "_execute_with_ptrs")` fast paths in experimental/ops/sdpa.py and the uid_order cache feeding them; and the five places docs/adding_torch_custom_ops.md told authors to hand-roll that path, which raises AttributeError as written. Backend execute on a 128^3 bf16 matmul: 16.17 -> 14.76 us. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * Fill a DLPack struct from a per-layout prototype, not field by field Every field of a `DLManagedTensor` except `data` is a property of the layout, yet `DeviceView.__dlpack__` built a fresh shape array and assigned nine ctypes fields on every call — and a graph makes ten of these per execute, one per workspace region carved for the kernel. Fill the struct once per (shape, dtype, device) and copy it: 1.68 us of field assignment becomes a 0.45 us memmove of 72 bytes. Measured 3.62 -> 1.39 us per `__dlpack__`, and GDN forward 154.5 -> 137.6 us end to end with no engine touched, because the ten workspace views are all it takes. This is the shape the backend already uses for kernel arguments (src/common/include/runtimeKernel.h): a prefilled blob plus, per mutable field, an (offset, uid, UpdateMethod) saying what execute writes where. Here there is exactly one mutable field, `data`, at a fixed offset, with update method POINTER, so the bookkeeping collapses to a memmove and one assignment. The struct stays FRESH per capsule. cute's from_dlpack aliases it rather than copying the DLTensor, so a struct shared between two capsules — or between two threads executing one graph — is read after someone else re-pointed it. Only the prototype is shared, and it is immutable. Also renames Operands to VariantPack: it IS the variant pack, normalized, and the python-side one being slightly wider than the C++ template's is not worth a second word. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * Drive the linear-attention engines from the normalized variant pack execute() already normalized the caller's operands once; the engines were still reading the caller's objects a second time. Every port went through buffers.probe(), which for bfloat16 falls out of __cuda_array_interface__ and into torch's __dlpack__ at 8.6 us apiece -- nine per GDN forward, 47.7 us, to learn dim and stride that the pack was already holding. _FrostPlan now takes the pack. The port-to-slot join is a property of the graph, so it is computed once and kept; between executes only the addresses move. What reaches the kernel is built from the pack rather than passed through, so the geometry a buffer is checked against and the geometry it runs on are the same reading. Contiguity moves with it, and becomes one gate instead of one call per compiled callable naming its own ports. That list was the same every time and had to be maintained by hand: a port added to a node but forgotten there went unchecked. Workspace joins too -- it needs a pointer, a device and a size, and the pack now carries all three, so Workspace.over() replaces a tenth probe. Two costs are added on purpose. Building eight DeviceViews is 10.1 us, and handing them to the kernel instead of the caller's tensors is another 17.6, because tvm-ffi reads a torch tensor through a C vtable (__dlpack_c_exchange_api__) and any python producer through a capsule. Both are the same fact -- python cannot build a fast DLPack producer -- and both go when the producer becomes a C type. Keeping the caller's tensor to avoid them would mean torch is a hard dependency of the engine path, which is the thing this removes. GDN forward, SM100, total=4096 H=4 D=128 4 seqs: before after contiguity gate 47.7 4.2 resolve_node_buffers 8.3 0 (bound once, kept) workspace probe 3.5 0.2 normalize 0 13.9 (now reads the workspace too) building the views 0 10.1 execute() 127 117 Also here, found while measuring: - selected_engine is a property execute() calls every time, and answering it walked every registered engine for the one declaring this id: 2.75 -> 0.48 us, cached against the plan config's identity so replanning invalidates it without a hook on every writer of _plan_index. - Two in-function imports of things the module already imports at the top. The one in VariantPack.view() ran per operand and cost 19 us of the GDN forward on its own. - _describe asked a torch tensor for its facts and gave everyone else a half-filled Tensor: no stride, no data_type. It now asks each producer in its own spelling -- torch's element-unit stride() and data_ptr(), cupy's byte-unit .strides and .data.ptr, one DLPack read for the rest -- so the same buffer is described the same way whoever produced it. fp8 is in the dtype table for the same reason; fp4 is deliberately not, since DTYPE_ITEMSIZE would make it zero bytes wide. - probe() declined two different ways through one exception, so an operand whose dtype has no name here lost its dim and stride as well. The two are told apart now. - The descriptor-skip cache in four kernels had a 0% hit rate: one of its guards compared against ws.view(...), a fresh object every call. It asked torch's _version counter whether cu_seqlens had changed, which was sound for torch callers and silently stale for everyone else. Deleting it is 5 us faster than keeping it. - check_buffer_device walked every operand asking which GPU it was on. cuDNN's own variant pack carries no device at all, and the question that matters is where the launch is going, not where the memory is: one current_device() read, 0.74 us against 1.45 per operand. 419 linear-attention tests pass, 1769 skipped. * Hold the variant pack as DLTensors, in one C type that is also the producer The pack was python objects: a Tensor per operand, and a DeviceView per operand and per workspace carve to hand the kernel. Both halves cost more than the work they describe. Reading a buffer meant asking a python object four questions one method call at a time and building a Tensor to hold the answers, 1.5 us each. Handing one back meant building a DLPack capsule, which tvm-ffi reads at 1.86 us where it reads a torch tensor at 0.35 -- through a C function table, __dlpack_c_exchange_api__, that no python producer can offer. Both halves are that one protocol, so this consumes it and implements it. VariantPackNative reads each operand through the caller's vtable into a DLTensor it keeps; the slots it hands out carry the vtable themselves, so a kernel reads ours through the same path it reads a framework tensor -- at 0.30 us, cheaper than the tensor it replaces. Refusing to pass the caller's object through therefore costs nothing, where insisting on it used to cost 17.6 us of kernel-argument conversion. A producer without the vtable is not an error and not a cliff: read_all returns the slots it could not take, python describes those with the reader it already had, and a mixed pack costs the sum of its parts. The workspace carves are the same type as the operands now, so a graph hands its kernels one kind of buffer rather than two, and DeviceView is off the hot path entirely. GDN forward, SM100, total=4096 H=4 D=128 4 seqs: before after normalize 13.9 6.3 contiguity gate 4.2 0.4 building the views 10.1 2.3 kernel-argument penalty 17.6 0 execute() 117 58 The backend path picks this up without a line changed: describe= had stopped selecting anything once reading was a single C call, so both paths take it and _execute_with_raw_ptrs reads the native pointer array directly. The parameter is gone. tensors[] is materialized on first access rather than built eagerly -- 16.9 us for eight operands, more than twice the whole normalize. Nothing on the per-execute path asks for it; frost_gemm will, for its M/N/K, and should read the native shapes instead when it migrates. No flag decides this: the laziness is the gate, and it needs no engine to declare anything. dlpack_version.txt moves 1.1 -> 1.3 for the DLPackExchangeAPI declarations. FetchContent keeps its checkout, so an incremental build needs _deps/dlpack-* cleared to actually pick the new tag up. Two things the migration surfaced, both in kernel code the forward path never reaches: - cute's from_dlpack at compile time does not read the vtable, so a slot needs __dlpack__ as well. It transfers ownership properly -- its own copy of the shape and stride plus a real deleter -- rather than aliasing storage the slot owns, which is how DeviceView's no-op deleter became a use-after-free whenever a consumer outlived the view. - the bprop state downcast reshapes its operand, so slots reshape too. A non-contiguous one is refused rather than silently reinterpreted: DeviceView could skip that check because it was row-major by construction, and a slot is whatever the caller passed. 430 linear-attention and dispatch tests pass, 1769 skipped. Two notes for anyone reading the numbers: - The fast path needs the producer's type to carry the vtable. torch 2.13 has it natively; on older torch tvm-ffi installs it by JIT-building a small extension, which is why flashinfer gets the same path there. So what decides it is whether tvm-ffi has been imported, not the torch version -- a backend-only process on old torch takes the python fallback, correctly and slowly. - The vtable is only called after walking prev_api for a table whose major version matches the header this was built against. The protocol requires that walk and keeps older tables reachable for it; without it a producer that moved to a new major version would have us calling function pointers at offsets it was free to move. * Initialize the dlpack dtype tables from either direction _FROST_DTYPE_CODE_TO_CUDNN was only ever filled by _dlpack_code_bits, which runs on the python fallback. An operand read through the exchange vtable never takes that path, so nothing populated the tables before something came looking for the reverse mapping and every VariantPack.tensors[i].data_type read back None -- silently, and only for the fast path. Verified before: BFLOAT16/FLOAT/INT32 operands all reported data_type=None. After: each reports its own. Nothing on the per-execute path reads this yet, which is why the suites stayed green; frost_gemm will when it migrates. Reported by coderabbit on #547. * Carve the workspace in one crossing, and settle the review comments The six regions a GDN forward carves are plan-time constants -- offsets fixed by WorkspaceLayout at build, dtypes and shapes fixed with them -- but every execute rebuilt them one at a time, at 0.9 us each: a bounds check, two dtype table lookups, an int() walk over the shape and a pybind crossing, per region. A carve compiled once at build hands back all six in one crossing (5.5 us to 0.8). GDN forward host time 58.5 to 51.0; the four frost engines all use it. Alongside, the review comments on the PR: - A python plan reached execute() with dynamic-shape overrides used to be handed the raw uid map, which a migrated plan cannot read. It cannot honour the overrides either -- a frost engine bakes the declared extents into the kernel it compiles -- so execute() refuses rather than answer a different problem than the caller asked. ExecutionContext's three override fields go with it: nothing could set them. - VariantPackSlot's DLTensor points into its own vectors, so its copy and move constructors are deleted rather than left to alias. - DeviceView.__dlpack__ handed out a struct it owned behind a no-op deleter, which a consumer outliving the view read after free. It delegates to a slot, whose capsule owns its struct and has a real deleter -- which retires the ctypes prototype machinery the view needed. - The exchange-vtable cache is keyed on a type's address, so it now holds a reference to it. - The reentrancy test gave eight threads one workspace to write. - L0 markers, CUDA gates, and two docs that described deleted code. * Normalize for a migrated plan whether or not overrides are passed The branch that skipped normalization when the caller passed override_uids / shapes / strides handed a migrated plan the raw uid map, which it cannot read. Refusing the overrides instead was wrong: frost_gemm compiles M/N/K symbolically and test_override_shape_frost runs other sizes through this exact call. They are accepted and change nothing here, so the branch goes. Also adds bench_sdpa_gemm_host.py, the counterpart of bench_gdn_host.py for the two engines that have not migrated: frost sdpa fwd 34.5 us, frost gemm 42.9, against GDN's 49.3 -- of which 14.8 is GDN's eight launches, so gemm carries the most host work of the three. * Migrate frost_gemm to the variant pack, and normalize in one crossing frost_gemm read the caller's buffer objects directly, which tied it to whatever framework produced them and -- because the graph declares B as [batch, K, N] while a caller allocates (batch, N, K) -- made it answer a different question than the backend under override_shapes. The two are one fix: the engine reads the pack, and execute() puts the overrides INTO the pack, so an engine honours them without knowing the concept exists. Overrides are re-expressed in the axis order the operand already uses. override_shapes speaks the graph's declaration; the slot holds what the caller's buffer reports. They are the same memory, so they rank their axes the same way by stride, and matching the two rankings gives the permutation. Applying the override verbatim left the pack describing the same bytes in a second language, and reading N off a fixed axis then read K. Test: same graph, same buffers, same override, backend and FROST both against the reference -- the case override shape is FOR, a max allocation with the live shape named per call. The existing coverage only checked FROST against itself, which is why this could diverge unnoticed. Normalization moved into the C pack while the engines were being pointed at it, since every path pays it: - read_from(uid_to_data, uids) does the lookups and the reads in one crossing, retiring the ordered list python built to hand to read_all - read_buffer_extent() reads the workspace through the same vtable; asking python for its size cost as much as reading all eight operands - first_unfilled() replaces a per-operand is_filled loop _normalize 6.0 -> 2.0 us. GDN forward 51.3 -> 45.3, frost gemm 42.9 -> 43.1 (the migration itself is free; what is left is normalization, which every engine now shares). run_resolved lets the engine skip rebuilding the by-object / by-uid / by-name tables on every execute. VariantPackSlot gains permute() and stride(dim) -- the kernel layer calls both on a caller buffer, and neither is visible from the engine directory. * Settle codex's findings, and prototype the gemm gate as one baked table Six findings from a codex review of the variant-pack work, five of them real and two memory-safety: - override_slot took ndim from the shape but stored whatever stride it was given, so a shorter stride array was read ndim deep by any consumer. It is the one place a shape and a stride arrive from two different lists; equal ranks are now required. - slot_managed_from_py_object shallow-copied the slot's DLTensor, whose shape and stride point into the slot's own vectors, with a deleter that freed only the wrapper. A managed tensor is the form a consumer may outlive the producer with, so it now owns copies. Same class of bug as the DeviceView no-op deleter fixed earlier -- the python half was repaired and the C vtable half was not. - read_buffer_extent computed a byte COUNT and Workspace.over read it as a byte RANGE, which is only the same for a dense buffer. A non-dense workspace is now refused rather than carved past its end. - A short override_shapes / override_strides silently kept the original metadata for the entries it did not name, where the backend rejects the request -- the same call, two geometries, decided by plan selection. - A bare device address normalized to a rank-zero unknown-dtype tensor, so an engine reading the pack for its extents failed on an operand form the backend has always accepted. It borrows the graph's declaration now. - The exchange-vtable cache kept null answers, but a type can acquire the vtable later (tvm-ffi installs one on import for torch builds without it), and a graph normalized before that import was pinned to the python fallback for the life of the process. Only hits are cached. The prototype, behind CUDNN_FRONTEND_FROST_GEMM_GATE_TABLE=1: which operand needs what alignment, what major, and how its extents follow M/N/K are settled when the plan compiles, so they are computed once into a table and walked once, instead of rebuilding five lists and walking them four times per execute. Measured on a 256x256x128 matmul: the gate alone 20.0 -> 13.3 us, graph.execute 43.7 -> 37.0. 4122 gemm tests pass with it on. The operands' pointer alignment stays with _alignment_reject rather than being re-checked inline: a gate emits a contract as well as a verdict, and a test matches its wording. One table evaluated twice keeps the message single too. Sizing it first was the point -- frost_gemm_execute_design.md makes the gate's share the acceptance condition, and it is 46% of execute. The same measurement names the next item: _call_positional is 14.1 us against a ~4 us floor of one launch plus one DSL crossing. watch_run.sh returns from a detached run on finished, KILLED or STALLED. The bare `until grep EXIT=` waiter only returns on the first, and a killed job then looks exactly like a running one. * Drop what nothing consumes, and stop explaining the obvious VariantPack.tensors materialized a Tensor record per operand for an engine that wanted the geometry as python objects. frost_gemm was that engine -- and when it was migrated it read the slots directly, which is cheaper and does not name a dtype the graph's vocabulary has to translate. So the property had no caller, and neither did the reverse dtype table built for it. A bug fixed in it earlier in this branch was a bug in code nobody ran, which is why it survived. read_all goes the same way: read_from does the lookups and the reads together, and nothing calls the older entry. The gemm gate-table prototype moves to its own branch. It is guarded by an env var and off by default, so in this PR it is a diff a reviewer has to read and cannot benefit from. The benchmark and watchdog scripts, and a design note for work that is not in this PR, come out of the tree entirely. Comments trimmed throughout: measurements and history belong in this description, not at a call site. What is left is the non-obvious and load-bearing -- why the override has to be re-expressed in the operand's axis order, why a byte count is not a byte range, why a managed tensor may not point at the slot's vectors. --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
Anerudhan
added a commit
that referenced
this pull request
Sep 2, 2026
* 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)
* remove unprofessional comments (#349)
* BSA: avoid guardword scanner false positives (#350)
* benchmark: fix repo-root path resolution in bench_moe (#348)
* Python-native cudnn.pygraph: graph IR + pluggable execution backends (#336)
* feat(python): backend-agnostic native graph + Router (unification proposal)
Modernize the Python-native graph API into the backend-dispatch architecture
from the Frontend v1 "Python API Engine and Graph API Unification" proposal.
Graph construction stays backend-agnostic; a backend is chosen by a first-class
Router at create_execution_plans() time (per Anerudhan's feedback), and the
backend-specific representation (e.g. the C++ cuDNN graph) is generated lazily
only then:
Python Graph API -> create_execution_plans() -> Router -> selected backend
(native engine, else cuDNN)
Layers kept separate:
- Graph IR (Node/Tensor/NativeGraph): engine-agnostic op DAG, full introspection
- BaseEngine: the backend contract (check_support/execute/get_workspace_size +
priority); cuDNN Graph is one routed backend, not a hardcoded default
- Router (engines/router.py): first-supporting by priority; None => cuDNN
Included: the IR, BaseEngine, Router, a CPU-only ReferenceMatmulEngine
(CI-testable correctness oracle), the optional MatmulCuTileEngine, and node
builders for block-scale / MoE / reduction so a DSL fusion backend can consume
them via graph.nodes (replacing the monkey-patch "recorder").
Deferred to follow-ups (see docs/python_native_graph_router.md):
NativeGraph.from_pygraph() (raises NotImplementedError for now), the DSL fusion
backend port, attention backends, and cuDNN lowering of the new node types.
Tests: 42 passing on CPU (IR + Router + reference-engine execute + cuDNN
fallback); cuTile path gated to SM100.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* refactor(python): trim NodeType to exercised ops; doc mixed candidate-list routing
- NodeType now lists only the op types this version exercises; drop the unused
norm/reshape/slice/etc. entries (re-add per-op when needed, following the
block-scale / MoE / reduction examples).
- Document the target routing model: create_execution_plans() takes one mixed
candidate list (native engines + cuDNN heur_modes) and produces a ranked list
of plans across backends; this PR ships the first-supporting-by-priority form.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* refactor(python): drop BATCHNORM / BATCHNORM_INFERENCE from NodeType (unused)
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* refactor(python): remove conv ops from native graph (unused foundation)
Drop CONV_FPROP / CONV_DGRAD / CONV_WGRAD: enum entries, the conv_fprop /
conv_dgrad builders, their dim inference in nodes.py, cuDNN lowering branches,
and the conv test. Re-add per-op when a backend needs conv, following the
block-scale / MoE / reduction examples.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs(python): use generic 'python DSLs' for backend examples
Avoid naming specific internal backends in public docs/docstrings.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* refactor(python): unify engines into one flat engine-id space (no cuDNN wrapper)
Replace the single-selected-backend + "if native else cpp" fork with the
engine-id model: python engines and cuDNN backend engines share one flat id
space. Python engines occupy a reserved high region (engine_ids.py,
PYTHON_ENGINE_ID_BASE = 1<<20) and each declares a stable engine_id it owns, so
ids never shift with registration order (reproducible autotune / pinned plans).
- engine_ids.py: PYTHON_ENGINE_ID_BASE + is_python_engine() + a phase-1
CUDNN_HEURISTIC_ENGINE_ID sentinel. Single source of truth for the namespace.
- Router.select()->one-engine becomes Router.plan()->ranked list of
PlanConfig(engine_id, knobs): supporting python engines (by id) + one trailing
cuDNN entry. TODO: interleave the true per-engine cuDNN configs
(get_engine_and_knobs_at_index) + real heuristics ranking; for now just concat.
- NativeGraph: _selected(engine) -> _plans(list) + _plan_index; add
get_execution_plan_count() / select_plan(i). check_support / build_plans /
get_workspace_size / execute all dispatch on the selected plan's id via
is_python_engine — one predicate, no fork. cuDNN is lowered lazily only when a
cuDNN-id plan is selected (pure-python when a python plan wins).
- BaseEngine: drop `priority`, add stable `engine_id` (reserved region).
reference_matmul = BASE+0, matmul_cutile = BASE+1.
Tests updated to assert the plan list; 41 pass on CPU incl. cuDNN fallback.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): make cudnn.pygraph engine-aware in place (transparent front door)
Users keep the classic API — g = cudnn.pygraph(...) is unchanged for every
existing sample — yet a graph transparently routes to a registered python engine
when it's fully represented. No new user-facing class, no rename.
pygraph_engines.install(pygraph) (called from __init__, same sanctioned pattern
as pygraph.execute = _execute) augments the pybind class in place:
- Per-graph mirror (WeakKeyDictionary) records a Node/Tensor IR alongside the
real C++ calls for a curated represented set (matmul + common pointwise),
mirrored via the NativeGraph builders so the recorded op is exactly what
engines consume.
- Every other op-builder is auto-wrapped to flag the graph "opaque" — the safe
direction: only disables the python path, never changes classic output.
- Lifecycle (create_execution_plans/check_support/build_plans/get_workspace_size/
execute/build) routes to a python engine iff one is registered AND the whole
graph is represented AND it supports the graph; else delegates to the untouched
C++ path.
Verified on an L40S against the real cuDNN build: a classic matmul runs
byte-identically with and without the augmentation, and a matmul+bias+relu graph
built via cudnn.pygraph + ReferenceMatmulEngine routes to the python engine with
exact results. Eager for now (C++ graph still built); lazy/pure-python is the
follow-up (needs a structured builder per op — multi-tensor returns like sdpa
can't be mirrored generically). NativeGraph stays as the standalone/greenfield
authoring object sharing the same IR + engines.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): native GEMM-family lowering + fix cuDNN execute path (phase 1)
Toward the native cudnn.pygraph migration (GEMM-family first). Make the native
build->lower->cuDNN execute path actually work end to end, and extend lowering
coverage to the GEMM family.
Fixes (all latent — the cuDNN execute path had never been GPU-tested):
- Thread the cuDNN handle: NativeGraph(handle=...) -> passed to the lowered
cudnn.pygraph so heuristics/build have a handle.
- Propagate the IR uid to the C++ tensor (was uid=-1 for auto tensors), so
execute()'s variant pack (keyed by IR uid) actually binds the buffers.
- POINTWISE lowering: the C++ pygraph has no generic pointwise(); dispatch on the
mode to the named ops (relu/gelu/sigmoid/tanh, add/mul/sub/div; add/mul also
cover bias/scale via broadcast).
Lowering coverage added: reduction, block_scale_dequantize, block_scale_quantize
(2 outputs), moe_grouped_matmul.
Validated on GPU (SM89): matmul and matmul+bias+relu built natively via
NativeGraph, lowered to cuDNN, execute with exact parity (new
test_native_cudnn_lowering.py, GPU-gated). Full native/router/pygraph suite: 45
passing. Per-op output-shape inference (e.g. reduction reduced dims) and
block-scale/moe execution parity are the next slices.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): reduction output-shape + SF reordering lowering (GEMM-family phase 2)
- reduction(): take an explicit reduced `dim` (cuDNN requires the reduction
output dims set); lowering sets set_dim/set_stride on the cuDNN op. Validated
matmul -> reduction(ADD over N) parity on GPU.
- lower_tensor(): propagate reordering_type to _make_tensor (e.g. F8_128x4),
needed for block-scale scale-factor tensors.
Native/router/pygraph + GPU parity suite: 46 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): native block-scale (nvfp4) lowering on Blackwell + fixes (phase 3)
Complete the GEMM-family native lowering with block-scale, validated on SM100.
Two more latent cuDNN-path bugs fixed:
- _lower_to_cpp passed io_data_type=None -> cudnn.pygraph rejects None. Now omit
io when unset; default intermediate/compute to FLOAT (matching cudnn.graph())
so cuDNN infers virtual (intermediate) tensor dtypes during build.
- lower_tensor now propagates reordering_type (F8_128x4) and omits data_type
when unset (NOT_SET) so cuDNN infers fused block-scale dequant output types.
Validated on SM100: dequant(A_fp4)@dequant(B_fp4) with F8_128x4 SFs builds +
executes via NativeGraph (test gated to SM100 + torch fp4; parity harness = the
repo's own fp4 test, which also only checks execution).
CPU overhead of the native Python layer (512^3 fp16, L40S): build +0.40 ms on
~106 ms (~0.4%, dominated by cuDNN heuristics); execute +0.3 us/call
(9.8 -> 10.1 us). Negligible.
Native/router/pygraph + GPU parity (matmul, bias+relu, reduction, block-scale):
48 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): native moe_grouped_matmul lowering + parity (GEMM-family complete)
- Add moe output-shape inference (token [1,T,H], weight [E,H,N] -> out [1,T,N])
so NativeGraph.validate() passes; cuDNN infers the same at build.
- GPU parity test (self-contained per-expert reference; no dependency on the
upstream test's helper) — validated on SM100.
GEMM family now fully native-lowered + validated on GPU: matmul, pointwise
(bias/relu), reduction, block-scale nvfp4, moe. Suite: 48 passing.
Next: non-GEMM ops (norms/reshape/slice/...) then the C++ _op rename + atomic
flip of cudnn.pygraph.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* fix(python): IR-uid -> C++-uid translation at execute; native rmsnorm (first norm)
Systemic fix: op-created C++ tensors (op outputs / virtuals) get uids assigned
by the C++ FE during build_operation_graph, in ITS enumeration order — which
does not match IR allocation order for multi-output ops (rmsnorm assigns
INV_VARIANCE=5, Y=6 while the IR allocated Y=5, inv_var=6). Keying the variant
pack by raw IR uids bound Y's buffer to inv_var: a [N,C,H,W] fp16 write into a
16-byte buffer (heap corruption / NaN). Single-output ops only worked by
allocation-order coincidence.
Fix: keep the lowering tensor_map; after build_operation_graph query every C++
tensor's real uid into an explicit IR-uid -> C++-uid map; execute() translates
variant-pack keys through it. No more order coincidence anywhere.
rmsnorm added as the first-class norm template (per "no corner-cutting" — the
generic opaque-op bridge was rejected/reverted since it makes non-GEMM ops
un-introspectable black boxes): named input/scale/epsilon/bias ports, Y/inv_var
outputs, norm_forward_phase param, pass-by-value epsilon; Y/inv_var dims carried
in the IR, cuDNN infers on its side. GPU parity: errY=0.0019, errI=0.0.
Suite: 49 passing (GEMM family re-validated through the translation path).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* refactor(python): Python IR owns the uid namespace end to end
Systematic uid review — four assignment paths existed:
1. user at creation: tensor(uid=...) (pybind _make_tensor, default -1)
2. user post-creation: tensor.set_uid() (mainline integrator pattern)
3. C++ FE auto-assign at build_operation_graph (enumeration order,
nondeterministic for multi-output ops) <- the coincidence trap
4. Python IR _alloc_uid (eager, sequential)
New invariant: for Python-built graphs, (3) NEVER triggers. The IR assigns
every uid eagerly at creation (auto or user-specified); lowering pushes ALL of
them explicitly to C++ — inputs via _make_tensor(uid=), op-created
outputs/virtuals via one set_uid loop over the complete tensor_map (single
point, impossible to forget per-op). Mixed construction (extending the lowered
C++ graph directly) is unsupported: a graph is pure-Python or pure-C++.
- Replace the IR->C++ uid translation map with a post-build ASSERTION: a
lowering path that fails to push a uid now fails loudly instead of being
silently translated (or worse, mis-binding buffers).
- _alloc_uid skips user-reserved uids; duplicate explicit uids rejected eagerly
at tensor() (C++ would only fail at build).
- execute() keys the variant pack by IR uids directly (== C++ uids by
construction).
Suite: 50 passing on SM100 (rmsnorm multi-output canary + block-scale included).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): full pointwise coverage — 54 ops, table-driven, mode == method name
Cover the entire pointwise surface of the C++ pygraph (54 methods) natively:
- Canonical op kind: params["mode"] IS the C++ pygraph method name (the
pointwise_mode enum is not exposed to Python; the method name is the semantic
name). Lowering collapses to a direct getattr dispatch — the mode<->method
mapping table is deleted as a concept.
- 47 uniform ops are generated from _POINTWISE_TENSOR_ARGS, a table of the
pybind tensor-argument names per op (mirrors the C++ signatures), so both
positional and the classic keyword call styles (bias(input=, bias=),
max(input0=, input1=)) work — required for the eventual cudnn.pygraph flip.
- 7 ops with scalar attributes get explicit builders storing them in params
(introspectable): relu(negative_slope/lower_clip/upper_clip), leaky_relu,
swish(swish_beta), gen_index(axis), + relu/leaky_relu/swish backwards.
Lowering forwards them as keywords.
- ReferenceMatmulEngine: keys move to method names; declines pointwise nodes
carrying scalar attributes it does not implement (correct-by-construction).
- Front-door mirror: classic calls passing scalar extras (e.g. relu clips) now
flag the graph opaque instead of silently dropping the attribute and
mis-routing to a python engine.
Tests: every builder exercised in both call styles + scalar-attr introspection
(CPU); sqrt/abs/max/min chain through real cuDNN on GPU. 53 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): norm family via one declarative table (10 ops, generic lowering)
All norms native — rmsnorm(_backward), layernorm(_backward), adalayernorm(_backward),
instancenorm(_backward), batchnorm, batchnorm_inference, batchnorm_backward —
through ONE mechanism instead of per-op code:
- _STRUCTURED_OPS: a declarative table per op — NodeType, tensor-input ports
(== the C++ pybind kwarg names), enum/scalar params (norm_forward_phase,
has_dbias), output ports in C++ return order, and per-output shape inference
(IR-side dims for introspection; cuDNN re-infers at build). Builders are
generated (keyword call style, as these ops are used repo-wide); lowering is
one generic branch: kwargs assembly + one call + zip outputs.
- List inputs (batchnorm peer_stats) become indexed ports (peer_stats_i) + a
count param, reassembled at lowering.
- The hand-written rmsnorm builder AND its lowering branch are deleted —
migrated into the table; the suite re-validates rmsnorm through the generic
path (multi-output uid canary intact).
GPU parity: layernorm fwd (Y/mean/inv_var) + layernorm_backward (DX/DScale/
DBias) vs torch autograd, using the supported LN config ([N,C,1,1]
channels_last, as in classic test_layernorm — the initial row-major 4D attempt
fails identically on the classic API, i.e. a kernel-support limit, not a
lowering bug). CPU: every table op builds a first-class node with named ports;
peer_stats port machinery covered. 56 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): conv + structural ops; collapse ALL structured ops into one table
_STRUCTURED_OPS now covers 25 ops — norms (11 incl. genstats), reduction,
block-scale (de)quantize, moe fwd/bwd, conv fprop/dgrad/wgrad, reshape, slice,
transpose, concatenate, rope fwd/bwd — one declarative entry each, one generic
lowering branch. Only matmul (positional ergonomics + front-door mirror) and
sdpa fwd/bwd (conditional kwarg assembly) remain explicit.
Deleted in the collapse: the hand-written reduction / block_scale_dequantize /
block_scale_quantize / moe_grouped_matmul builders AND their four lowering
branches, plus nodes.py moe shape inference (moved to the table). The suite
re-validates all of them through the generic path on GPU.
Table mechanics extended (each a one-word spec key, no new concepts):
- attrs: scalar/enum/list params forwarded verbatim (padding vectors, axis,
slices, permutation, reshape_mode, rope_dim, mode, ...). Conv accepts BOTH
the symmetric `padding` convenience and pre/post_padding — forwarded as
given; pybind overload resolution picks the right C++ binding.
- out_dims reserved kwarg (list, or {port: dims}): explicit output shapes for
ops cuDNN cannot infer — generalizes reduction's old `dim` param.
- push_output_dims: IR dims pushed to C++ for dgrad/wgrad/reduction/reshape/
moe_bwd (classic API also requires set_dim there).
- no_cdt: bindings without compute_data_type (reshape, concatenate).
- Builders accept tensors positionally or by port name; infer lambdas are
best-effort (try/except -> None; C++ validates at build).
GPU parity added: conv_fprop vs torch conv2d (NHWC), incl. asserting the
table's shape inference. CPU: all 25 ops x 2 call styles + out_dims. 58 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python): sdpa family via generic kwarg capture — full ~130-arg surface
The six sdpa variants (sdpa, sdpa_backward, sdpa_fp8, sdpa_fp8_backward,
sdpa_mxfp8, sdpa_mxfp8_backward) are now declared in _CAPTURED_OPS, the third
and final table mechanism: builders capture ALL kwargs generically — tensor
values (incl. torch/dlpack) become named ports (port == C++ kwarg), scalars /
enums / score_mod callbacks go to params verbatim, dropout tuples are flattened
per element — and lowering rebuilds the kwargs for one C++ call. The full C++
kwarg surface (~130 args: paged attention tables, diagonal bands, sink tokens,
cu_seqlens, fp8 descales/amaxes, ...) is supported without hand-mirroring any
of it, and future binding args are picked up automatically.
Deleted: the explicit sdpa/sdpa_backward builders (~170 lines, common-args
only) + their two lowering branches + nodes.py sdpa shape inference (moved to
table lambdas — and fixed: O is q-shaped with v's head dim, not v-shaped).
Semantics now match the classic API exactly: sdpa always returns (O, Stats)
with Stats None in inference mode (generate_stats/is_inference logic); output
dim/stride are pushed to C++ (the SDPA node requires O's layout pre-validate —
that's how BSHD vs BHSD output is chosen).
GPU: sdpa causal fp16 EXECUTION parity vs torch SDPA (was build-only before).
59 passing.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* feat(python)!: THE FLIP — cudnn.pygraph is now the Python graph class
The public cudnn.pygraph name now binds the Python IR class (class name:
pygraph; module: python/cudnn/pygraph.py — no "relative-to-history" naming).
The C++ graph builder is internal-only at cudnn._pybind_module.pygraph and is
reached exclusively through lowering: a graph is pure-Python or pure-C++,
never mixed. Zero C++ changes — the demotion is by namespace, not rebuild.
Deleted in the flip (afterthought residue):
- pygraph_engines.py front-door + its tests (no install()/monkey-patching
anywhere: register_backend is a native method on the class)
- NativeGraph.from_pygraph stub (meaningless now), use_native back-door
- docs/python_native_graph_router.md (initial-brainstorm doc, per review)
Drop-in surface for classic parity, driven by iterating the repo's own test
files until green (each item below was a real failure caught and fixed):
- conditional outputs ("maybe"): rmsnorm_backward(has_dbias=False) -> DBias
None; norm fwd INFERENCE -> mean/inv_var None; batchnorm next_running_*
present iff in_running_* given (classic returns None for absent outputs)
- torch interop: tensor(dim=x.size()) (torch.Size), data_type=torch.bfloat16
(converted at the C++ boundary via _library_type, IR stores user's value)
- output dtype semantics: an output without explicit set_data_type gets io
dtype (was mis-defaulted to intermediate FLOAT -> fp32 into fp16 buffers)
- Tensor gains the classic setter/getter surface (set_ragged_offset,
set_reordering_type, set_is_pass_by_value, ...); tensor_like(cudnn tensor);
tensor_scalar; CPU tensor_like -> pass-by-value (classic rule)
- ragged (THD) output layout: outputs' ragged_offset now pushed to C++ at all
mapping sites (was silently dense -> wrong values in sdpa_thd)
- validate-time table shape inference (topological): chained ops whose inputs
are virtual (conv on a relu output) infer once inputs are known;
builder-time infer stays as best-effort for direct inputs
- classic lifecycle: build_operation_graph lowers eagerly when no python
engines are registered, so deselect_*/query methods work between classic
steps via __getattr__ delegation to the lowered graph; build_plans(policy)
passthrough; deserialize(*args, **kwargs) passthrough incl.
enforce_precompiled; execute override_uids/shapes/strides + dlpack pointers;
get_execution_plan_count = python engines + backend's dynamically-queried
count (frontend NEVER statically enumerates backend engines — they vary by
backend version; Router keeps ONE delegating cuDNN entry by design)
- stride optional after set_dim (row-major inferred), None variant-pack keys
tolerated, C++-tensor keys resolved via get_uid
Validated: our suite (56) + classic spot-runs all green on real GPUs —
matmul_bias_relu, rmsnorm, layernorm, batchnorm, conv_fprop (incl.
execute_plan_at_index), apply_rope, kernel_cache, sdpa_with_caching, sdpa_thd,
sdpa_chunked_prefill (ragged+paged), conv_genstats, conv_reduction, slice,
block_scale_quantize_dynamic_shape, wgrads. Full-suite runs on SM100 + mhas in
flight; residuals to follow. Known pre-existing env skew (fails identically on
the unflipped installed package): test_deviceless_aot_compilation on this box.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): classic validate() timing + omit unset compute_data_type
Two classic-parity fixes surfaced by the full mhas run (3567 uniform failures,
one root cause):
- cudnnGraphNotSupportedError must fire at graph.validate(): the classic test
waiver pattern is try/except-skip AROUND validate(), with
build_operation_graph() called bare. With no python engines registered,
validate() now lowers and runs the C++ validate right there (unsupported
configs skip, not fail); build_operation_graph()/plan creation are staged
behind flags so each C++ step runs exactly once in classic sequencing.
Python-engine graphs still never touch C++ at validate.
- compute_data_type=None is now OMITTED at every lowering site (matmul /
pointwise / structured / captured) instead of passed through: classic ops
default to NOT_SET in C++; pybind rejects None. Also converts via
_library_type when set (torch dtype parity).
Previously-failing mhas case now skips as on classic; our suite 56 passing.
Full-suite + full-mhas reruns in flight.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* docs(router): codify the extension contract for the future heuristics MR
Ranking policy is intentionally undecided; what IS decided: policy pluggable at
three levels (Router subclass / per-graph / process default); plan() may return
any ordering or mix; backend engine sets are discovered per graph at plan time
(never statically enumerated); PlanConfig can carry concrete backend engine
configs, with pygraph._lower_cudnn_plan as the designated point to honor them.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): plan-selection lifecycle + registration validation (review items 2, 6)
Review item 2 (reproduced bugs):
- ONE plan index space: [0, n_python) are python plans, [n_python, ...) are the
backend's plans (sub-index = index - n_python, queried dynamically).
get_execution_plan_count() and select_plan() now agree; selecting a backend
sub-index lowers on demand, builds via build_plan_at_index and executes via
_execute_plan_at_index (sub-index 0 == the classic default path).
- select_plan() survives build()/execute(): build() no longer silently re-plans
when a plan list exists (explicit create_execution_plans() still re-plans).
Review item 6:
- register_backend() validates at registration: engine_id must be a stable int
in the reserved python region, unique per graph; registration after planning
is rejected. BaseEngine.engine_id defaults to None so a subclass that forgets
to declare identity fails clearly instead of silently colliding.
- Decline signal narrowed: an engine declines ONLY via NotImplementedError or
cudnn.cudnnGraphNotSupportedError (the classic unsupported-graph signal);
ValueError/RuntimeError now propagate as engine bugs instead of silently
falling back to cuDNN. Reference/cuTile engines updated accordingly.
Regression tests for all of the above (pin-survives-execute, duplicate/missing
id, post-planning registration, unexpected-exception propagation). 59 passing +
classic spot files green.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* feat(python): compiled-plan engine lifecycle + ExecutionContext (review item 1)
The engine contract now represents a real JIT/DSL backend:
- propose_plans(graph) -> [PlanConfig]: one engine may expose several
configurations to ranking/autotune (default: one plan with default_knobs
when check_support accepts). PlanConfig moves to engines/base.py.
- build_plan(graph, plan) -> CompiledPlan: the expensive JIT step, run ONCE per
(graph, selected plan) at build_plans() time. The compiled artifact is cached
ON THE GRAPH (keyed by plan index), so one engine instance is safely reusable
across graphs and repeated execution reuses the artifact. The selected plan's
knobs reach build_plan verbatim.
- CompiledPlan.get_workspace_size(): plan-specific workspace; graph
get_workspace_size() reports it for python plans.
- ExecutionContext(handle, stream, workspace, override_uids/shapes/strides)
passed to CompiledPlan.execute(): stream resolved from the caller's handle
(classic cudnn.set_stream semantics); caller workspace object reaches the
plan; no engine hard-codes a stream (cuTile now launches on ctx.stream).
- Simple eager engines are unchanged in spirit: implement execute() only; the
default build_plan wraps it in a trivial CompiledPlan.
Acceptance tests per the review: two knob proposals from one engine with the
selected plan's knobs observed at build+execute; compile-once artifact reuse
across executions; same engine instance on two graphs without state collision;
plan-specific nonzero workspace; caller workspace object identity at execute.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): IR port direction, tensor identity ownership, parity gaps (review items 3, 4, 5, 7)
Item 3 — SDPA capture direction:
- _CAPTURED_OPS entries declare out_kwargs (rng_dump, score_max, score_sum_exp,
dBias, dSink_token): tensor kwargs that are semantically OUTPUTS are recorded
in node.outputs (correct producer/consumer for engines) and still forwarded
as descriptor args at lowering. fp8/fp8_backward positional schemas extended
to the full binding order (descales/scales).
Item 4 — tensor identity is graph-owned:
- Tensor hash/eq are object identity (uid/name are mutable; value hashing broke
the dict-key invariant). set_name/set_uid delegate to the owning graph
(weakref set at registration) which re-indexes atomically: name index, uid
index, auto-bound data follow; duplicate names and USER-user uid conflicts
raise. Classic-parity subtlety the review didn't cover: classic tensors have
no uid until set_uid while the IR assigns eagerly — a user set_uid landing on
an auto-assigned uid silently renumbers the auto holder (auto uids are
internal until lowering) instead of failing classic code.
Item 5 — parity gaps: get_workspace_size(*args) classic overload passthrough;
serialize() lowers on demand (cuDNN-format by definition, independent of the
selected plan); stale references to the removed design doc dropped.
Item 7 — freeze policy: structural mutation (new ops via the _get_name
chokepoint, tensor rename/re-uid, backend registration) raises after
lowering/planning instead of desynchronizing derived state.
Classic gaps found by the SM100 full-suite sweep (fixed + re-validated):
- slice: classic passes `slices` POSITIONALLY -> structured builders now map
extra positionals onto attrs in declared order (covers conv paddings too);
output dims inferred from the python slice objects; output dtype inherits the
input's (dtype_like), matching the C++ rule.
- moe_grouped_matmul: token_index/token_ks ports + top_k attr (gather/scatter).
Environment skew documented (fails identically on the unflipped installed
package; installed .so older than repo tests): test_mhas_v2 sdpa_mxfp8
(`implementation=` kwarg not in installed binding) and
test_deviceless_aot_compilation (`enforce_precompiled`).
122 tests green locally (contract + classic spot files incl. set_uid-heavy
kernel-cache/sdpa-caching).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): address coderabbit inline findings (broadcast checks, cuTile hardening, tensor_scalar parity)
- matmul batch broadcast: incompatible extents raise (numpy rules) instead of
silently taking max.
- pointwise broadcast inference: right-aligned merge across ALL inputs;
lower-rank operands no longer dropped; incompatible extents raise.
- MatmulCuTileEngine: CUDA runtime return codes checked (failures decline the
engine); execute verifies all operands share one CUDA device (multi-GPU
hosts: mismatched context silently corrupts).
- tensor_scalar: scalar_type is required (classic binding takes it positionally
in every overload) — also closes the lowering path where an untyped
pass-by-value scalar silently dropped its embedded value.
Two other findings were already fixed before these comments were filed:
default engine_id collision (registration validation, BaseEngine.engine_id =
None) and mutable-uid Tensor hashing (identity hash/eq).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): review follow-up — replan invalidation, slot-based dispatch, context/freeze/validation completeness
Follow-up item 1 (stale artifact on explicit replan): create_execution_plans()
now invalidates every plan-derived artifact (compiled python plans, built
state, the backend's plan list) — a stale compilation can never execute.
Follow-up item 2 (mixed Router ordering): dispatch is slot-based, honoring the
Router's ordering verbatim. _plan_slots() maps every public index to
("python", PlanConfig) or ("cudnn", sub_index) with the cuDNN entry expanding
in place; selection, workspace, build and execute all use the same mapping.
cuDNN-first and interleaved orderings now work as the router contract promises
(prefix-count assumptions removed).
Follow-up item 3 (context completeness): build_plan(graph, plan, ctx) receives
a build context (handle + stream) — no private-state reads for AoT compilers.
Stream resolution is strict: a supplied handle whose stream query fails RAISES
(never a silent stream-0 fallback); with no handle, engines resolve
deterministically from their framework (cuTile: torch current stream).
Dynamic workspace-query overrides on python plans are rejected explicitly
instead of silently ignored.
Follow-up item 4 (MXFP8 schemas): match the bindings exactly — full positional
orders (fwd: +descale_q/k/v; bwd: q_T/k_T/o_f16/dO_f16/dO_T + all descales),
dSink_token as an output kwarg, named outputs (dQ,dK,dV,amax_*); rng_dump
removed from fp8_backward (not on that binding).
Follow-up item 5 (freeze completeness): ALL semantic Tensor setters (dim,
stride, data_type, output/virtual, ragged, reordering, pass-by-value) are
frozen after lowering/planning via the owner guard; tensor_scalar registers
through _register_tensor (owner installed, identity mutations re-index).
Follow-up item 6 (validation bypasses): constructor-provided backends go
through register_backend() validation; propose_plans() results are checked for
foreign engine-id injection; duplicate explicit tensor names are rejected at
initial registration; CUDA runtime API failures in cuTile propagate as
RuntimeError (an unsupported arch/driver remains a normal decline).
Acceptance tests for each item (replan invalidation, interleaved-router
dispatch, constructor/proposal validation, workspace-override rejection,
strict stream failure, mxfp8 port direction, per-setter freeze, scalar
ownership, duplicate names). 74 contract tests + classic spot files green.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* refactor(python): one-shot planning (classic conformance) + retire NativeGraph name
Planning is one-shot: a second create_execution_plans() raises. Empirically the
classic C++ graph never supported re-planning (a second call there APPENDS
plans by accident, build_operation_graph twice hard-errors, and mutation after
build is silently stale) and no user re-plans. The replan-invalidation
machinery added for review follow-up item 1 defended a capability that had no
users — deleted; the same guarantee (a stale compiled artifact can never
execute) now holds structurally because plan state is write-once. Autotune
re-selects WITHIN one plan set via select_plan(), matching the classic
build_plan_at_index flow. Plan differently => build a new graph (IR
construction costs microseconds).
Also retire the transitional NativeGraph name everywhere (tests, engine
docstrings, type hints) — the class is cudnn.pygraph, full stop. A single
documented alias line remains for downstream migration.
109 tests green.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): stable two-level plan indices; land the two missed patches (review round 3)
Round-3 review items:
1. STABLE plan indices (the lazy-expansion contradiction): the flat in-place
expansion of the cuDNN entry shifted python plans' indices when lowering
happened (index 2 became a cuDNN sub-plan, python-B moved to 4) — pinning
was unreliable. Adopted the two-level model the original review sanctioned:
top level = the Router's entries verbatim (each python PlanConfig one index,
the cuDNN delegating entry ONE stable index = the classic default path);
backend sub-plans stay in the backend's own index space via the classic
build_plan_at_index / execute_plan_at_index / *_plan_at_index APIs
(delegated). Indices never shift; the expansion machinery is deleted.
get_execution_plan_count keeps the exact classic semantic when no python
engines are registered.
2. C++ replan-appends: moot since planning became one-shot (83ffdedcf) — the
C++ create_execution_plans can no longer be reached twice on one graph
(enqueue_engine_configs appending was exactly why replan had to go).
3. Landed for real (previous patches missed their anchor strings and failed
silently — now grep-verified): cuTile resolves torch's current stream when
no handle stream exists (literal stream 0 gone); rng_dump removed from the
fp8_backward schema (not on that binding). Also: execute()-supplied handle
now reaches the JIT build on auto-build (the python path plans first and
compiles with the caller's ExecutionContext instead of running the generic
build with only the graph handle).
4. Custom-Router bypass closed: create_execution_plans() validates the FINAL
router output — python entries must name registered engines, only one cuDNN
delegating entry allowed, anything else raises.
5. get_dim()/get_stride() return copies (the classic pybind getters return
fresh lists; live-list mutation after planning is no longer possible).
111 tests green.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix: callback graph shim for score_mod closures; serialize returns classic form
Two classic-parity fixes found by running the full suite on a current
extension build:
- flexible SDPA score_mod callbacks: user closures capture IR Tensors but
the callback receives the lowered C++ graph. _CallbackGraphShim translates
IR Tensor arguments at the call site (lowering closure-captured helper
tensors on demand), so existing callback code runs unchanged.
- serialize(): return the C++ binding's serialized form unchanged instead
of wrapping in bytes. C++ deserialize casts the payload back to
vector<uint8_t> and rejects bytes, so the bytes wrapper broke the classic
serialize -> deserialize(handle, data, enforce_precompiled=True) round
trip (test_deviceless_aot_compilation::test_device_properties).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): review round 4 — explicit planning state, split plan-index spaces
- get_execution_plan_count() is ALWAYS the classic backend-count passthrough
(lowering the cuDNN entry on demand); it never returns the routed-list
length, so its semantics no longer depend on whether python engines are
registered. The routed plan list is graph.plans / select_plan() — a
separate, stable index space. An unplanned graph counts 0 (classic), and a
python-only routed graph raises with a pointer to graph.plans.
- Explicit _planning_done flag replaces the nonempty-list proxy everywhere
(one-shot check, register_backend, set_router, freeze, build/execute
needs-planning checks); an empty Router output is rejected — there is no
legal empty planning state. set_router after planning raises.
- cuTile resolves the fallback stream on the OPERANDS' device
(current_stream(a.device)), after the same-device check — argless
current_stream() is the active device's stream, which can be a different
GPU on multi-GPU hosts.
- router.py contract downgraded to what this MR enforces: at most one cuDNN
delegating sentinel; concrete cuDNN engine configs as routed entries are
the heuristics follow-up's typed-plan work, not one extra lowering branch.
- tensor(uid=) creation path now applies the same collision rule as
set_uid: a user uid landing on an auto-assigned uid steals it (holder
renumbered); only user-user collisions raise. Found by the SM100
block_scale_quantize dynamic-shape tests, which assign explicit uids after
ops already auto-assigned.
Tests: cuDNN slot of a mixed router actually executes through the backend
with routed indices stable across lowering (GPU); one-shot planning on a
pure-cuDNN graph (GPU); empty router rejected; set_router frozen after
planning; backend-count/routed-space separation; creation-path uid steal.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): push ragged_offset_multiplier on output tensors at lowering
The three output-mapping sites pushed set_ragged_offset but not the
multiplier, so a non-default multiplier on an output (unified SDPA ragged
layouts, paged fp8 fwd) lowered as multiplier=1 — the backend computed wrong
addresses (cudaErrorMisalignedAddress, hard process abort). The input path
already passed it via _make_tensor kwargs.
Found by full test_mhas_v2 -m '' on H100/dev-9.26: 21x
test_sdpa_random_fwd_ragged_offset_multiplier_unified_L1 + 1x
test_sdpa_fp8_fwd_paged_L0 crashed on the flip and passed on the
classic-control package (same .so, develop python files). After the fix the
same selection is 145 passed, matching classic exactly.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): push reordering_type on output tensors; consolidate output-attr lowering
Same bug class as the ragged multiplier: an attribute set on an OP OUTPUT via
the classic setter chain (block_scale.set_reordering_type(F8_128x4) in
test_block_scale_quantize) was never pushed at the output-mapping sites, so
the backend rejected the quantize scale layout on SM100. The three duplicated
output blocks are consolidated into one push_output_attrs helper (ragged
offset + multiplier, reordering, output flag, dtype) so the next
output-settable attribute has exactly one place to go.
Attribution: 7 test_block_scale_quantize failures on Blackwell were
flip-attributable (classic control passes); fixed. The 3 test_cudnn_sdpa_op
d=256 failures fail identically on the classic control (environment).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): whole-surface freeze + output layout contract; split cuTile engine out
Review round 5:
- Freeze covers the ENTIRE public surface, not just the fluent API. An
explicit _frozen flag is set at lowering and at planning (whichever
first); _freeze() seals node port/param dicts to MappingProxy views,
dim/stride lists to tuples, and Tensor/Node/GraphContext gain __setattr__
guards. graph.nodes / graph.tensors return copies. A mutation while
merely validated (python-engine graphs stay mutable until planning)
invalidates _is_validated so stale inference never reaches planning.
- Output layout contract: Tensor tracks user-assigned vs IR-inferred
dim/stride; push_output_attrs pushes USER-assigned layouts verbatim
(previously lost on matmul/pointwise outputs) and never pushes inferred
row-major strides — the backend keeps its classic per-op inference
(channels-last conv). Tests: explicit column-major matmul output stride
honored end to end; conv output stays channels-last in the lowered JSON.
- cuTile matmul engine split out of this PR (engine file, optional extra,
tests, exports) — it re-lands with the DSL-engine integration PR;
ReferenceMatmulEngine remains the in-tree contract oracle. This PR is the
contract, not a kernel product.
- MoE lowering test gated on cuDNN 9.15+.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* refactor(python): retire pygraph name collision; drop NativeGraph; check in design doc
Review feedback:
- The C++ pybind graph class is renamed pygraph -> backend_graph
(cudnn._compiled_module.backend_graph): two things named pygraph was
confusing now that cudnn.pygraph IS the Python class. Internal-only
rename — nothing public imported the pybind name post-flip.
- The Python module moves to cudnn/_pygraph.py (private module, public
class re-export), so the class qualname is cudnn._pygraph.pygraph, not
the double-take cudnn.pygraph.pygraph.
- NativeGraph transitional alias dropped completely.
- Design doc checked in: docs/python_graph_and_execution_backends.md —
architecture, two plan-index spaces, engine contract, invariants
(uid ownership, one-shot planning, freeze, output layout), naming, and
follow-up scope.
- test_native_cudnn_lowering: every cuDNN-path execute now asserts
dispatch-level proof it ran through the backend plan path
(_assert_ran_on_cudnn: cuDNN entry selected, graph lowered, backend
plans created/built). Kernel identity below the backend API is
deliberately not asserted — kernel names are backend-internal and
version-dependent; numerics + dispatch proof is the stable contract.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* refactor(python): 'cudnn' never means 'the backend' in names — both sides are cuDNN
The frontend Python graph is as much cuDNN as the C++ library; identifiers
that used 'cudnn' to designate the backend side now say 'backend':
- CUDNN_HEURISTIC_ENGINE_ID -> BACKEND_HEURISTIC_ENGINE_ID
- _lower_cudnn_plan / _has_cudnn_plan / _cudnn_heuristics ->
_lower_backend_plan / _has_backend_plan / _backend_heuristics
- _assert_ran_on_cudnn -> _assert_ran_on_backend
- test_native_cudnn_lowering.py -> test_native_backend_lowering.py
(tests *_lowers_to_cudnn -> *_lowers_to_backend, mixed-router /
one-shot test names likewise)
- docstrings/comments: 'cuDNN entry/sentinel/slot/path/side' ->
'backend ...' throughout; 'the cuDNN C++ backend' stays where it
describes what the backend is.
Also fixes a stale TYPE_CHECKING import left by the module rename.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): classic-parity batch from internal CI — signatures, wrapper, labels, naming, layout truth
Root-caused from the internal CI failures (py_samples / pycudnnTest); every
item below reproduces 1:1 against the classic package on the same GPU/backend
and is fixed + validated (pycudnnTest 26/26, all 13 CI sample notebooks pass,
local battery 2172/0):
- Constructor and tensor() are POSITIONALLY IDENTICAL to the classic API
(name is the constructor's first positional arg — pycudnnTest passes it
positionally; classic sm_count/sm_version/kernel_cache/device_property/
dynamic-shape params explicit; classic tensor() order with is_pass_by_value/
ragged_offset/reordering before name/uid; NOT_SET/-1/NONE sentinels
normalized). New params (backends/router) are keyword-only. Guarded by
test_api_signature_parity, which reads the classic order from the pybind
docstring/wrapper itself.
- wrapper.py (cudnn.Graph) recognizes IR tensors: one _GRAPH_TENSOR_TYPES
tuple replaces 7 isinstance(cudnn.tensor) sites (the notebooks' silent
UnboundLocalError/mis-capture).
- Duplicate tensor names are legal classic LABELS (pycudnnTest builds two
'weight's): uid is identity; the name index serves unique names only and
ambiguous-name lookups raise instead of guessing.
- Op outputs are auto-named with the classic C++ conventions
(node::MEAN/INV_VARIANCE/DSCALE..., per-op overrides for rmsnorm_backward's
::Dscale/::Dbias) — wrapper.Graph canonical-name lookups depend on them.
- Multi-output ops return a LIST like classic pybind (pycudnnTest dispatches
on isinstance(res, list)).
- Layout truth: backend-inferred dim/stride are reflected back into the IR
after build_operation_graph (_sync_ir_shapes_from_backend) — wrapper
allocates output buffers from IR getters; provisional row-major strides are
no longer observable post-build. push_output_dims ops push stride only when
USER-assigned (pushing inferred row-major into an NHWC graph made the
backend reject dgrad+add fusion).
- tensor_like normalizes non-torch DLPack objects (CuPy .strides is in
BYTES) through torch.from_dlpack — NHWC CuPy inputs no longer silently
become row-major.
- get_data_type() returns the cudnn enum when the user stored a torch dtype
(classic converts at set time).
- validate() no longer auto-marks leaf outputs as non-virtual — discarding a
result (training SDPA's Stats in the paged sample) is legal classic usage;
auto-marking made its uid required in the variant pack.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* chore: remove internal test file accidentally included
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* fix(python): renames are label writes (exempt from freeze); push output names at lowering
Two classic-parity items from the internal CI notebook set:
- set_name after build is legal classic usage (sample 24 renames a tensor on
an already-built graph): names are labels with no execution semantics, so
_rename_tensor no longer consults the freeze — the label write bypasses the
sealed-tensor guard explicitly, and the ambiguity policy still governs the
name index.
- User renames on op OUTPUTS now reach the lowered graph: push_output_attrs
pushes the IR name, matching classic where the rename acts on the same
object the cpp graph holds (visible in JSON dumps and wrapper.Graph
canonical-name lookups).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* test: skip introspection/one-shot tests when cudnn.pygraph is monkey-patched
The internal tree layers a DSL engine by monkey-patching cudnn.pygraph
lifecycle methods process-wide at import (cudnn.TBD). Under pytest-xdist any
worker that collects those tests carries the patches into unrelated tests:
signature introspection then sees the wrapper's (*args, **kwargs) and the
patched create_execution_plans swallows the one-shot error (except Exception)
— false negatives against pristine-class contracts.
Detect the replacement via __qualname__ and skip LOUDLY with the reason,
instead of failing on behavior that is not this class's. The proper fix
remains scoping the internal patches (fixture install/uninstall) or excluding
the TBD shard from the shared py_test run; these guards just make the
contamination visible as skips rather than red.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---------
Co-authored-by: Yang Xu <yanxu@nvidia.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* Add NWH + B2B causal conv1d notebooks; refresh outputs (#246)
* Add NWH + B2B causal conv1d notebooks; refresh outputs
- Add 62_causal_conv1d_nwh_forward.ipynb
- Add 63_causal_conv1d_nwh_backward.ipynb
- Add 64_b2b_causal_conv1d_forward.ipynb
- Add 65_b2b_causal_conv1d_backward.ipynb
- Refresh outputs for all 6 notebooks (60-65)
* Guard NWH and B2B causal conv1d APIs for cuDNN 9.24
* Format causal conv1d Python op
* Match CI Black line length for causal conv1d op
* Add runtime guard for causal conv1d 9.24 symbols
* Address CodeRabbit B2B causal conv1d feedback
* Document causal conv1d notebook version requirements
* Allow zero grad for discarded B2B output
---------
Co-authored-by: Hwanseo Choi <hwanseoc@nvidia.com>
* Bump development version to 1.27.0 (#358)
* Fix FE-OSS docs links and DSA architecture code fence. (#360)
Use stable SDPA documentation URLs in overview and mark the DSA architecture block as text to avoid code highlighter parsing issues.
* Fix cutlass DSL deprecation: use .ptr for cute.struct scalar fields (#365)
cutlass-dsl 4.5+ deprecates using a @cute.struct scalar field directly
as a pointer (_ScalarData.value), emitting:
DeprecationWarning: Use explicit `struct.scalar.ptr` for pointer instead.
from cute/core.py whenever tmem_holding_buf / tmem_dealloc_mbar_ptr are
passed to cute.arch.alloc_tmem / retrieve_tmem_ptr / utils.TmemAllocator.
Switch the remaining call sites to the explicit .ptr accessor, matching
the pattern already used by the other DSA kernels.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* Fix DSA offset alignment, stream handling, and CUDA Graph capture (#354)
* Fix SM90 query offset alignment
Signed-off-by: kunlunl <kunlunl@nvidia.com>
* Preserve the CUDA default stream
Signed-off-by: kunlunl <kunlunl@nvidia.com>
* Make dense indexer backward graph safe
Signed-off-by: kunlunl <kunlunl@nvidia.com>
* Require a tensor grad_loss for indexer backward
Signed-off-by: kunlunl <kunlunl@nvidia.com>
* Address remaining DSA review comments
Signed-off-by: kunlunl <kunlunl@nvidia.com>
---------
Signed-off-by: kunlunl <kunlunl@nvidia.com>
* Expose cu_seq_len_q/kv on the sdpa_fp8 python binding (#366)
* Expose cu_seq_len_q/kv on the sdpa_fp8 python binding
The unified-engine FP8/MXFP8 forward (cuDNN 9.25+) accepts cumulative
sequence lengths, and the C++ API has supported them on the fp8 node since
1.25 (SDPA_fp8_attributes aliases SDPA_attributes), but the python sdpa_fp8
binding hardcoded cu_seq_len_q/kv to nullptr. Expose them as kwargs
(appended last to preserve positional backward compatibility) so python
callers can use fp8 + cu_seq_len; the python-native pygraph capture/replay
layer forwards them without changes.
- python/pygraph/{pygraph.h,sdpa.cpp}: add cu_seq_len_q/kv parameters to
PyGraph::sdpa_fp8 and its m.def, with docstring entries (requires cuDNN
9.25+ and the UNIFIED implementation). Remove a stale "Deprecated, use
sdpa_unified instead" comment: implementation selection is automatic (or
explicit via the implementation attribute), and python fp8 users are
expected to call sdpa_fp8.
- docs/operations/Attention.md: document cu_seq_len_q/kv on the fp16/bf16
C++ and python APIs (missed in #266), the ragged offset multiplier
(missed in #290), and the fp8 varlen surface incl. the new kwargs.
- test/python/sdpa/fp8.py: support is_cu_seq_len and
with_ragged_offset_multiplier configs (mirroring fp16.py): cu_seq_len
graph tensors, token-coarse offsets with per-tensor multipliers on
Q/K/V/O, version gating at 9.25.
- test/python/test_mhas_v2.py: test_sdpa_fp8_fwd_ragged_L0 now draws
ragged / cu_ragged / cu_ragged_mult with equal weight.
sdpa_mxfp8 is intentionally untouched: it has no varlen surface at all
(no padding mask or seq_len kwargs), so cu_seq_len support there is a
separate feature.
Validated against cuDNN 9.25 (test_sdpa_fp8_fwd_ragged_L0): H100 10
passed / 22 skipped (pre-existing Hopper config limits), Blackwell 24
passed / 8 skipped (head-dim limits); the passing draws include 23
is_cu_seq_len=True and 9 multiplier configs, zero failures.
* Complete cu_seq_len docstring constraints on sdpa_fp8
Address review: the runtime-visible docstring now carries the same
constraints as the sdpa() docstring and Attention.md — set together,
use_padding_mask=True, cuDNN 9.25+ and the UNIFIED implementation.
* Serialize selected plan behavior notes (#364)
* Serialize selected plan behavior notes
* Add behavior note serialization regression sample
* Add pip install --group dev, prerequisite for deprecating requirements.txt (#359)
* Add dev dependency group
* Reorder pyproject sections
* Fix uncaught ValueError in flatten_pass_by_value on malformed hex input (#343)
The hex branch of flatten_pass_by_value converted "0x"-prefixed strings
without error handling, so malformed values such as "0x" or "0xZZ" in a
log's pass_by_value field crashed the cudnn_repro CLI with an unhandled
ValueError. Guard the conversion with the same try/except pattern the
decimal branch already uses, returning an empty list for unparseable
strings, and add regression tests.
Fixes https://github.com/NVIDIA/cudnn-frontend/issues/342
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* Make plan structure serialization optional within serialize() to construct symmetry with deserialize logic (#371)
* Make plan structure serialization optional within serialize() to construct symmetry with deserialize logic
* add CUDNN_FRONTEND_UNUSED for guarded out macro case
* Update SDPA Benchmarking Artifacts - 9.24.0.43 (#362)
* Organize FE OSS tests by feature (#372)
Organize FE OSS tests into flat feature directories and update imports and documentation paths.
* GEMM+RoPE+MXFP8 fusion (#367)
* Add fused gemm+rope+mxfp8quant kernel.
* Add documentation and general interface names
* Address coderabbitai's suggestions
* Additional tests for fused gemm+rope+mxfp8
* Remove NUM_HEADS constant
* Test: organize GEMM projection tests (#374)
* Fix cutlass DSL deprecation warnings in CuTe DSL kernels (#376)
Fixes all cutlass-dsl 4.5.x deprecation and optimization warnings
emitted by the CuTe DSL kernels during the OSS test suite:
- tcgen05.OperandMajorMode -> cute.nvgpu.OperandMajorMode (also
silences the <string>:11 warnings raised inside the MMA op ctor
when the deprecated enum type is passed through).
- make_trivial_tiled_mma / make_blockscaled_trivial_tiled_mma legacy
single-ab_dtype overload -> new overload with separate a_dtype and
b_dtype (dtype duplicated, matching the legacy path exactly).
- cutlass.utils.distributed.atomicAdd -> local dsl_user_op wrapper
over cute.arch.atomic_add with identical relaxed/sys semantics.
- Static loops with >=64 iterations flagged by DSLOptimizationWarning:
cutlass.range_constexpr -> cutlass.range(..., unroll_full=True)
where the loop body only needs dynamic tensor indexing.
No functional changes; codegen is equivalent.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* Fix BSA backward hang on cute-dsl 4.6.0: version-gate elect_one around bulk stats copies (#382)
cute-dsl 4.6.0 changed cute.copy lowering for bulk-async atoms
(cpasync.CopyBulkG2SOp, TMA): the copy now elects a single lane
internally via a warp-collective WARPSYNC.COLLECTIVE + ELECT.
bsa_bwd_sm100's load warp wrapped its LSE/dPsum stats copies
(cute.copy with CopyBulkG2SOp) in cute.arch.elect_one(), as required
on <= 4.5.x where the bulk copy did not self-elect. On 4.6.0 the two
elects nest: lane 0, alone inside the outer elect region, reaches the
copy's internal warp-collective elect which waits for all 32 lanes and
deadlocks the warp. Q/LSE/dO/dPsum stop flowing and every other warp
spins in mbarrier waits; in CI the oss:rel [Blackwell] job pegged the
GPU at 100% until the 1h job timeout
(https://gitlab-master.nvidia.com/cudnn/cudnn_frontend/-/jobs/360031309).
Diagnosed by cuda-gdb break-in on the live hang (2 TMA-load warps
parked at WARPSYNC.COLLECTIVE/ELECT inside the stats copy; 26 warps
spinning in SYNCS.PHASECHK downstream) and by PTX A/B diff showing
stacked double elect.sync at the stats-copy sites on 4.6.0 vs a single
one on 4.5.0.
Fix: introduce copy_utils.bulk_copy_elect_one(), which returns
cute.arch.elect_one() on cute-dsl <= 4.5.x and a nullcontext on
>= 4.6.0, and use it at the four copy_stats sites. All other
elect_one uses (mbarrier init/arrive, consumer_release, tcgen05
commits, cp.reduce.async.bulk inline asm) still require the guard and
are unchanged.
Verified on Blackwell (SM 10.0):
- cutlass-dsl 4.6.0: test/python/fe_api/block_sparse_attention
17 passed in 31.7s (previously 3 device-side hangs)
- cutlass-dsl 4.5.0: unchanged behavior via the version gate
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* Fix architecture-independent SDPA repro failures (#386)
* Restore SDPA repro tensor dumps
* Skip FP8 reference checks in perf mode
* Fix large MXFP8 performance repros
* Address SDPA repro review comments
* Keep MXFP8 storage access direct
* Use UID map for tensor dump collection
* Avoid monkeypatching MXFP8 performance test
* Use vector for tensor dump collection
* Remove MXFP8 performance smoke test
* Add SDPA edge case test coverage (#328)
* Add SDPA edge case tests
* Refine SDPA edge case coverage
* Add cu_seqlen zero-length edge tests
* Guard cu_seqlen tests by cuDNN version
* Run zero seqlen tests from cuDNN 9.25
* Improve GitHub issue and pull request templates (#375)
* Infra: improve GitHub issue and PR templates
* Infra: make issue forms less restrictive
* Infra: simplify bug environment fields
* Infra: simplify feature request form
* Infra: expand bug environment prompt
* Infra: combine CUDA environment versions
* Infra: clarify optional bug environment
* Infra: remove redundant GPU environment field
* Infra: add cuDNN version examples
* Infra: consolidate bug and PR templates
* Infra: limit pre-commit reminder to staged files
* Infra: format PR area choices vertically
* Infra: simplify CodeRabbit auto-review config
* Support SM90 DSA qh16 indexer forward and fix qh32 sparse backward (#388)
* Support SM90 DSA qh16 and fix sparse backward
Addresses NVIDIA/cudnn-frontend#373 and NVIDIA/cudnn-frontend#385.
* docs: correct DSA SM90 support overview
---------
Co-authored-by: mingyangw <mingyangw@nvidia.com>
* Remove dead BSA fragment allocations (#392)
* Add collect_env environment report tool for bug reports (#400)
Issue reporters often can't state their environment precisely, and the
most common unreproducible-issue root cause is version confusion:
multiple cuDNN/CUDA copies installed where the loaded one is not the
one the user assumes.
python -m cudnn.collect_env produces an offline, read-only report:
frontend/backend versions with mismatch flags (stale pip metadata,
torch's libcudnn vs the frontend's dlopen'ed backend), the frontend's
libcudnn search-order resolution, GPUs in CUDA enumeration order,
loaded-vs-on-disk GPU libraries via /proc/self/maps with pip
provenance, relevant packages incl. torch's declared cuDNN pin, and
CUDNN_*/CUDA_* env vars.
Stdlib-only at module level with every probe individually guarded, so
the file also runs standalone with bare Python when import cudnn is
broken. Referenced from the bug-report issue template and README.
Co-authored-by: Yang Xu <yanxu@nvidia.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* Update nvidia-cutlass-dsl version to 4.6.0 (#368)
* Update nvidia-cutlass-dsl version to 4.6.0
* Migrate warp redux to public cute.arch.warp_redux_sync for cutlass-dsl 4.6.0
nvidia-cutlass-dsl 4.6.0 renamed the nvvm dialect enum ReduxKind to
ReductionKind, breaking every kernel that imported it and failing 284
Blackwell OSS tests at import time.
Instead of chasing the private-API rename, drop the three repo-local
redux helpers (moe_kernel_helpers.warp_redux_sync,
discrete_kernel_utils.warp_redux_sync, utils.warp_redux_sync_fmax, and
rmsnorm's redux_sync_max_f32) and call the public
cute.arch.warp_redux_sync(value, kind="fmax", ...) wrapper everywhere,
matching the pattern already used by gemm_srelu/gemm_dsrelu and the
DSA kernels. Also replace the raw nvvm.redux_sync bitcast sequence in
gemm_amax, whose res= kwarg was likewise removed in 4.6.0.
The old local helpers hardcoded redux.sync.max.abs.NaN.f32 regardless
of argumen…
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
cudnn.pygraphis now a Python-native graph class: graph structure (nodes, tensors, params) lives in Python with full introspection, and execution dispatches through pluggable backends — python DSL engines and the cuDNN C++ backend. The routed plan list is python plans + at most one backend delegating entry (the backend's own plans stay behind it in the classic at-index space); concrete backend engine configs as first-class routed entries are the heuristics follow-up MR's job (typed plan representation), not this one. The C++ graph builder is internal-only (cudnn._pybind_module.backend_graph), reached exclusively through lowering. The only C++ change is the pybind class renamepygraph->backend_graph(internal-only; nothing public imports the pybind name post-flip). Classic-API compatibility is validated empirically against the repo's own test suite on real GPUs (see Validation); known remaining parity items (constructor/tensor positional signature order, an API-inventory test, classic Tensor alignment/vector queries) are tracked explicitly and listed under Follow-ups.Implements the Python API Engine and Graph API Unification proposal (router at
create_execution_plans, lazy lowering, backend-agnostic construction).Architecture
graph_types.Tensor,nodes.Node,pygraph.pygraph: engine-agnostic op DAG. Input/output port names == the C++ pybind kwarg names, everywhere.engines.BaseEngine, a plan→compile→execute lifecycle:propose_plans(graph)(several knob configs per engine),build_plan(graph, plan, ctx) → CompiledPlan(JIT once per graph/plan, cached on the graph),CompiledPlan.execute(graph, tensor_data, ExecutionContext)with explicit handle/stream/workspace/overrides. Each engine owns a stableengine_idin a reserved region (PYTHON_ENGINE_ID_BASE = 1<<20). Simple eager engines implementexecute()only. Engines consumegraph.nodesdirectly.PlanConfig(engine_id, knobs)list in ANY ordering/mix (python-first, backend-first, interleaved); the Router's final output is validated regardless of Router implementation (registered ids only, at most one backend sentinel). Plan indices are two-level and stable: top level = the Router's entries verbatim (each python PlanConfig one index, the backend delegating entry ONE index — indices never shift on lowering;select_plan()operates here); backend level = the backend's own plans via the classic at-index APIs (build_plan_at_index/execute_plan_at_index), delegated verbatim. Ranking policy is intentionally undecided — the extension contract for a future heuristics MR is codified inengines/router.py(policy pluggable at 3 levels; backend engine sets discovered per graph at plan time, never statically enumerated — the frontend must work against any backend version).Key invariants
get_execution_plan_count()— lowers the backend entry on demand; that is the caller asking for the backend.)create_execution_plans()raises — classic never supported replan (empirically it appends engine configs on re-call and hard-errors on a secondbuild_operation_graph); switch plans viaselect_plan(), re-plan by building a new graph. After planning, structural mutation (new ops, tensor renames, semantic setters) raises.cudnnGraphNotSupportedErroratvalidate(); outputs withoutset_data_typeget io dtype; conditional outputs returnNone(e.g.has_dbias=False, INFERENCE-phase stats); torch dtypes/torch.Sizeaccepted and converted at the C++ boundary; ragged (THD) offsets on outputs;deserialize/build_plans/execute-overrides passthrough; plan-config/query methods delegate to the lowered graph.Op coverage (100% of the C++ op surface, three declarative mechanisms)
_POINTWISE_TENSOR_ARGSmode== method name, lowering is onegetattr_STRUCTURED_OPS_CAPTURED_OPSValidation
execute_plan_at_index), apply_rope, kernel_cache, sdpa_with_caching, sdpa_thd, sdpa_chunked_prefill (ragged + paged), conv_genstats, conv_reduction, slice, wgrads.Follow-ups (separate MRs)
Routerpolicy + a typed plan representation so concrete backend engine configs (engine_id, knobs) can appear as first-class routed entries next to python plans (contract boundaries inengines/router.py).BaseEngine(its probe/build registry, plan-list semantics and pinning match this architecture); its lifecycle monkey-patches and op-recorder become deletions (~1000 lines) since the graph is natively introspectable. The cuTile matmul engine was split out of this PR into that track —ReferenceMatmulEngineremains as the in-tree contract oracle._is_validated/_is_built/_planning_done/_cpp_*), aCudnnBackendAdapterto remove theselected_engine is Nonespecial-casing, extract lowering into its own module, dedupe op identity (NodeType enum vs registry keys), longer-term a typedOpSpecas the single per-op source for builder/validation/lowering.tensor()positional-signature order matched to the classic pybind signatures + a public-API inventory test (old vs newpygraph/tensor surface);🤖 Generated with Claude Code
Summary by CodeRabbit