Make plan structure serialization optional within serialize() to construct symmetry with deserialize logic - #371
Merged
hwanseoc merged 2 commits intoJul 9, 2026
Conversation
…truct symmetry with deserialize logic
Contributor
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (1)
🚧 Files skipped from review as they are similar to previous changes (1)
📝 WalkthroughWalkthroughGraph::serialize now accepts an optional ChangesOptional structural serialization
Estimated code review effort: 3 (Moderate) | ~20 minutes 🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
Contributor
Author
|
@cudnn-ci-bot run |
|
🚀 Running mirror pipeline Branch: cudnn-gh/pr-371-9a9ca76 |
Contributor
Author
|
@cudnn-ci-bot run |
|
🚀 Running mirror pipeline Branch: cudnn-gh/pr-371-aa5562f |
Anerudhan
added
mod-frontend
hwanseoc
approved these changes
Jul 9, 2026
Member
There was a problem hiding this comment.
Approved. This works standalone and the serialization/repro tests pass. A later PR currently assumes graph JSON v2 structural fields (json_version, gid, and tensors) are present during handle-based deserialization; that PR will be adjusted after this merges to explicitly support the plan-only payload produced by serialize_structure=false.
Merged
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
4 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
This PR makes plan serialization symmetric with the handle-based deserialize path and adds regression coverage for both payload modes.
Changes
Reason
Previously, serialization always emitted structure and always collected pass-by-value/workspace data from the node subtree. That could degrade plan-only round trips after deserialize(handle, ...), where structural graph nodes may not be present; particularly in cases where
deserialize(const json&, ...)isn't called. This change preserves correctness for plan-reload workflows while keeping default behavior backward-compatible.Summary by CodeRabbit