Make a family a kind of graph, and stop paying for the DSL to decline

[YangXu1990uiuc]

Before submitting

• I agree to license this contribution under the terms of LICENSE.txt.
• I ran pre-commit run and committed any formatting changes.
• I added GitHub labels: cat-cleanup, mod-frontend, mod-frost, orig-nv-eng.

Affected area

Python API or bindings

Summary

Cleanup of the python engine dispatch that landed with #476. Each change deletes a declaration that duplicated a judgment made elsewhere — and every one of those duplicates was already wrong.

An sm100 kernel now says it serves the sm100 line

Capabilities.arches was an exact set, frozenset({(10, 0), (10, 3)}), checked with device_cc not in arches. Rubin (sm107) and Thor (sm110) are meant to reuse these kernels and both were excluded. It is a range now — SM100-119 — and the decline message says so instead of listing device families.

The manifest carried a second, narrower copy, which is deleted:

family	manifest said	the implementation says
frost_gemm	SM100-103	PIPELINE_ARCH_RANGES["sm100"] = ((100, 120),), commented "family-portable Blackwell instructions"
frost_sdpa_bwd	SM120-121	Capabilities 120-129

On Rubin the gemm family would not have been offered at all, so its engine never got asked. Whether an engine suits a device is said once now, by the engine.

A family is a kind of graph, not a group of engines

Eight manifest entries were two different things wearing one name: three owned a 100-wide id block and returned several engines, five owned a single id and returned one. GDN and KDA had two entries each declaring identical anchors, so those graphs were claimed by two families.

Six families now, and classification is a lookup (_ANCHOR_NODE_TO_FAMILY) rather than N families each declaring a claim that then has to be proven disjoint. family_for(graph) is a pure property of the graph — no sm, no environment — because what kind of graph something is cannot depend on which machine is asking. Availability is a separate question.

The manifest assigns engine ids

Ids were assigned in four places (two _ID_OFFSETS tables, two class attributes) and the manifest only validated containment afterwards. instantiate() hands each factory the ids its engines are to use; engines declare none. opt_in moves with them, per engine — the half-precision SDPA engines can graduate while the fp8 one matures.

register_backend() is now only the out-of-tree escape hatch its docstring always claimed. Eleven tests called it on in-tree engines, all redundant since the manifest discovers them; with them go the "in-tree owner may register itself" exemption and its test.

closed_under is deleted, and the RESHAPE bug is fixed at its source

It promised gemm served RESHAPE when nothing did, so matmul -> reshape was claimed and then failed in execute demanding a buffer the caller never bound. 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. It declines now.

Deciding costs no framework

That deletion is only safe because declining stopped costing an import. Three package __init__s are lazy (PEP 562) and the DSL adapter resolves at build time:

	before	after
import cudnn.sdpa.graph_analyzer	1059 ms, +381 modules, pulls cutlass	9.4 ms, +2 modules
support-check module	+1387 modules	+7 modules

Facts and capabilities speak cudnn.data_type rather than torch.dtype, and the device comes from cudnn.create_device_properties() — the backend's own descriptor, the same object the C++ deviceless-AoT path serializes. graph_analyzer imports no torch at all now. test_import_boundaries.py holds the line, in a fresh interpreter, measuring the delta against an empty one.

A missing or too-old CuTe DSL is a decline at check_support, via probes that never execute the module (find_spec 7 ms, importlib.metadata 5 ms, against ~4.3 s for the real import). ImportError also joins decline_types() as a second line of defence — without it, deferring the import had turned "extra not installed" from "the family is absent" into an exception that skipped past the backend sitting in the same ranked list.

Why

Every declaration this deletes duplicated a judgment made elsewhere, and every one of the duplicates was already wrong: closed_under promised RESHAPE support nothing implemented, the manifest's arch range capped frost_gemm at SM103 against templates declaring [100, 120), and Capabilities.arches enumerated device families rather than the line, excluding Rubin and Thor. The shared shape is allow-by-default over an open set — a coarse copy that says yes where the real check would say no.

Related issues

Follows #476 (FROST engines). Related to #506, which fixes a CUDA-graph-capture bug in buffers.probe() that this investigation surfaced.

API and compatibility impact

None for callers. EngineFamily, the manifest, _attach_facts / _facts_for are internal; lazy exports (from cudnn.sdpa import SdpafwdSm100D256) are preserved by PEP 562. Python engine ids are renumbered, which engine_ids.py documents as permitted while pre-release — no API could return one before create_execution_plan() learned about them.

EngineRow is gone rather than deprecated: it arrived with #476 two days before this PR, has never been in a release, and carried anchors/closed_under, both deleted by design — a shim preserving them would be a new place to lie.

pyproject's cutedsl extra stays at >=4.5.0. The engines need 4.7.0, but pinning that here would make cudnn-frontend incompatible with anything holding the DSL back — quack-kernels pins ==4.6.0. The floor lives at runtime instead, where it costs only the engines that need it: check_support() declines an older DSL and the graph goes to the backend, and the test gate uses the same predicate so the suite skips rather than fails.

Behavioural: matmul -> reshape declines at build instead of failing inside execute. That path could not have worked before.

Testing

Blackwell sm100:

sdpa/frost + gemm/frost + test_mhas_v2  6678 passed, 2826 skipped   (-n 16)
test/python/test_engine_router.py         66 passed
test/python/test_import_boundaries.py      5 passed

The 4500/2119 split is identical to the counts before this series began.

test_mhas_v2 matters separately: it is the only suite that goes through routing rather than pinning an engine and exercising the kernel, and it reports which engine served each graph. FROST serves 413/3201 graphs across four engines (d128 182, d256 122, d128_fp8 55, d192_d128 54), so consolidating families, moving id assignment into the manifest, widening the arch declaration and deferring the DSL leave routing intact.

It also caught a bug the other suites structurally could not: deferring torch had put the import inside one branch of the execute path while a later, independent branch used it too, so an fp8 graph hit UnboundLocalError. Thirty tests. sdpa/frost stayed green throughout because it never reaches that path.

On the public wheel, which nothing had been verified against — this box ships nvidia-cutlass-dsl-internal while pyproject declares nvidia-cutlass-dsl. Swapped internal 0.3.0 for public 4.7.0: 4500 passed / 2119 skipped, identical.

Across architectures, which matters because deleting the manifest's arch filter opened a path sm100 cannot exercise — the engines are now imported and must decline themselves:

sm80 (A100)   sm89 (L40S)   sm90 (H100)     engines decline cleanly, no DSL pulled

While there: twelve gemm registry tests had no arch gate at all and fail on any non-sm100 host on develop today; they now carry requires_sm100. The SDPA suite's five hand-copied _is_sm100 gates, each pinned to exactly (10, 0) and so skipping on sm103, are one marker aligned with what the engines declare.

test/python/linear_attention fails 286 / passes 67 here — the same 286/67/1769 on pristine gh/develop, verified in a separate worktree. CUDA graph capture, untouched by this PR.

Design doc

docs/python_graph_and_execution_backends.md is updated in this PR — it has been checked in since #336 but described the shape this replaces. Four sections added: what the manifest decides and what it leaves to the engine, how facts are attached and why after the freeze, how handle/stream/device resolve, and which imports each dispatch stage may pay for.

Known gaps

• create_execution_plan(engine_id, knobs), the explicit replay entry point, does not go through finalize/freeze/attach; only create_execution_plans() does.
• A custom Router calling key() can retry backend lowering after facts were attached.
• _freeze() is shallow: MappingProxyType(dict(node.params)) freezes the mapping, not list-valued params.
• A failed planning attempt leaves the graph frozen while register_backend() / set_router() still gate on _planning_done, so an engine registered afterwards is ignored by the cached candidate set.
• heuristics_sort still does not read facts. Deliberate, and its docstring says what will: order a family's engines on its facts, then merge against the backend on predicted time.

[coderabbitai]

Note

Reviews paused

It 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 reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

• @coderabbitai resume to resume automatic reviews.
• @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

• ▶️ Resume reviews
• 🔍 Trigger review

📝 Walkthrough

Walkthrough

The router now classifies graphs by engine family, finalizes backend layouts, freezes graphs, and attaches analyzer facts before ranking. SDPA analysis uses cuDNN data types and cached facts. SDPA and FROST exports load lazily.

Changes

Engine family routing

Layer / File(s)	Summary
Engine family manifest and classification python/cudnn/engines/manifest.py, python/cudnn/engines/__init__.py, python/cudnn/engines/engine_ids.py, python/cudnn/engines/heuristics.py	Manifest rows are replaced by EngineFamily. Graphs resolve to one offered family, analyzers resolve lazily, and engines receive validated family IDs.
Family engine factories and identifiers python/cudnn/linear_attention/..., python/cudnn/gemm/frost/engine.py, python/cudnn/sdpa/fwd/engine.py, python/cudnn/sdpa/bwd/engine.py	Factories accept manifest-provided IDs. Linear-attention and FROST families use dedicated identifier blocks.
Backend planning and analyzer facts python/cudnn/_pygraph.py, python/cudnn/gemm/frost/graph_analyzer.py, python/cudnn/frost/buffers.py, python/cudnn/frost/README.md	Planning finalizes layouts, handles backend declines, freezes graphs, attaches facts, and ranks Python plans before backend plans.
SDPA facts and capability checks python/cudnn/sdpa/graph_analyzer.py, python/cudnn/sdpa/fwd/engines.py, python/cudnn/sdpa/bwd/engines.py, test/python/sdpa/frost/test_sdpa_graph_analyzer.py	SDPA facts use cuDNN data types and record sequence-length inputs. Engines consume cached facts and resolve DSL, CUDA, and Torch dependencies lazily.
Lazy exports and import boundaries python/cudnn/sdpa/**/__init__.py, python/cudnn/linear_attention/frost/__init__.py, test/python/test_import_boundaries.py	Public symbols resolve on first access. Import-boundary tests verify that graph analysis and support checks avoid heavy dependencies.
Routing and integration validation test/python/test_engine_router.py, test/python/linear_attention/frost/*, test/python/gemm/frost/test_frontend_integration.py	Tests cover family classification, ID allocation, analyzer caching, manifest consistency, and default routing without manual backend registration.

Estimated code review effort: 4 (Complex) | ~60 minutes

Sequence Diagram(s)

sequenceDiagram
  participant Graph
  participant BackendLayout
  participant Manifest
  participant Analyzer
  participant Router
  Graph->>BackendLayout: finalize layout
  BackendLayout->>Graph: freeze graph or record decline
  Graph->>Manifest: resolve family
  Manifest->>Analyzer: resolve and analyze family graph
  Analyzer->>Graph: attach cached facts
  Graph->>Router: generate ranked plan

Loading

Suggested reviewers: anerudhan

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 57.55% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.

✅ Passed checks (4 passed)

Check name	Status	Explanation
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Description check	✅ Passed	The description completes all required sections and provides detailed scope, rationale, compatibility impact, related issues, testing, and known gaps.
Title check	✅ Passed	The title clearly summarizes the main changes: graph-based family dispatch and deferred DSL imports for clean engine declines.

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

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

[YangXu1990uiuc]

Earlier comments on this PR were deleted rather than left to mislead: they described the pre-rescope design, and one claim in them is now backwards — I wrote that fp8 SDPA could not be split into its own family because its ids had shipped. engine_ids.py states the opposite (pre-release, nothing persisted them), which is why this version re-cuts the blocks freely.

Two things from those comments are worth keeping, both about measuring these suites:

These suites randomize shapes from an unseeded random. s_q, s_kv, d_qk, d_v, h_k/h_v and the mask bounds all come from random.choice; @torch_fork_set_rng(seed=0) forks only torch's RNG, i.e. tensor values. Two runs therefore compare different graphs — I measured roughly 49 tests of drift per run before noticing. Any before/after comparison on test_mhas* needs random.seed() pinned first, e.g. via a -p plugin outside the tree so both sides load the identical one.

-k is a substring match. -k "test1 or test2 or test3" selects test1, test10–test19, test100–test199, … — 182 tests, not 3. I misread my own filter and published a "cost is severely sublinear, so compilation dominates" conclusion from it. The real cost is close to linear at ~1.45 s/test with ~75 s of one-time per-process warmup, and the parallelism figures in the description supersede the ones I gave then.

[coderabbitai]

Actionable comments posted: 6

🧹 Nitpick comments (3)

python/cudnn/_pygraph.py (2)

1013-1037: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value

Align the _attach_facts docstring with the single-family lookup.

The docstring says the graph "finds its own families" and "hangs one record per family off itself". manifest.family_for() returns at most one family, so the code attaches at most one record. The plural wording will mislead a reader into expecting multi-family attachment.

📝 Proposed doc fix

-        Part of planning, not something a caller invokes: the graph finds its
-        own families through the manifest and hangs one record per family off
-        itself, as an optional payload. A graph no family claims carries none.
+        Part of planning, not something a caller invokes: the graph finds its
+        own family through the manifest and hangs that family's record off
+        itself, as an optional payload. A graph no family claims carries none.

🤖 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/_pygraph.py` around lines 1013 - 1037, The _attach_facts
docstring describes multi-family discovery and attachment, but
manifest.family_for() returns only one family. Update the docstring to use
singular wording throughout, stating that the graph finds its family and
attaches at most one family-scoped record, while preserving the existing
no-family behavior.

---

1002-1011: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Extract the backend-decline rollback into one helper.

Lines 1004-1011 and lines 1129-1142 record the same decline and roll back the same five fields. The two sites must stay in step; if one gains a field and the other does not, a half-lowered graph survives and a later build_operation_graph() walks into the descriptor that already failed. Extract a _record_backend_decline(exc) helper and call it from both.

♻️ Proposed refactor

+    def _record_backend_decline(self, exc: Exception) -> None:
+        """Record a backend decline and roll back any partial lowering.
+
+        A half-lowered graph makes a later build_operation_graph() walk into
+        the descriptor that just failed."""
+        self._backend_declined = exc
+        self._lowered_graph = None
+        self._cpp_tensors.clear()
+        self._cpp_bog_done = False
+        self._cpp_plans_created = False
+        self._backend_entries = []
+
     def _finalize_backend_layout(self) -> None:

Then in _finalize_backend_layout:

         except (cudnn.cudnnGraphNotSupportedError, RuntimeError, ImportError, AttributeError) as exc:
             _LOG.warning("backend could not build this graph, treating as a decline: %s", exc)
-            self._backend_declined = exc
-            self._lowered_graph = None
-            self._cpp_tensors.clear()
-            self._cpp_bog_done = False
-            self._cpp_plans_created = False
-            self._backend_entries = []
+            self._record_backend_decline(exc)

🤖 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/_pygraph.py` around lines 1002 - 1011, Extract the duplicated
backend-decline state reset into a `_record_backend_decline(exc)` helper,
including logging the exception, storing `_backend_declined`, clearing
`_lowered_graph` and `_cpp_tensors`, resetting `_cpp_bog_done` and
`_cpp_plans_created`, and emptying `_backend_entries`. Replace the rollback
blocks in the `_lower_backend_graph()` exception path and
`_finalize_backend_layout` with calls to this helper so both paths remain
synchronized.

test/python/test_engine_router.py (1)

877-878: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Point the probe family at a real factory to avoid a cached failure.

"unused_factory" does not exist in this module. manifest.instantiate() therefore raises AttributeError, logs a WARNING with a full traceback on every passing run, and caches [] under engine id _OOT + 900 in the module-level _INSTANCES. monkeypatch restores MANIFEST but not _INSTANCES, so that entry outlives both tests. A later test that reuses id _OOT + 900 with a working factory would receive the cached empty list instead.

Define a module-level factory that returns [] and name it here.

♻️ Proposed test fix

+def _probe_factory():
+    """No engines: the probe family exists to carry an analyzer, not a kernel."""
+    return []
+
+
 def _probe_analyzer(graph):

Then at both construction sites:

-    family = manifest.EngineFamily(_OOT + 900, "probe_family", __name__, "unused_factory", analyzer=(__name__, "_probe_analyzer"))
+    family = manifest.EngineFamily(_OOT + 900, "probe_family", __name__, "_probe_factory", analyzer=(__name__, "_probe_analyzer"))

Also applies to: 920-921

🤖 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_engine_router.py` around lines 877 - 878, Define a
module-level probe factory in test_engine_router.py that returns an empty list,
then replace the invalid "unused_factory" value with that factory’s name in both
EngineFamily construction sites for IDs _OOT + 900 and _OOT + 920/921. Keep the
existing analyzer configuration unchanged.

🤖 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/_pygraph.py`:
- Line 151: Update the inline comment on the _facts field to describe that it
stores analyzer records only and reference the _facts_for() accessor instead of
facts(). Do not change the _facts data structure or surrounding logic.

In `@python/cudnn/engines/__init__.py`:
- Line 28: Restore the EngineRow compatibility export in cudnn.engines,
preserving its legacy constructor and behavior through a deprecated adapter;
update the package exports such as __all__ and imports alongside MANIFEST and
EngineFamily so existing cudnn.engines.EngineRow imports continue working.

In `@python/cudnn/engines/manifest.py`:
- Around line 24-32: Update the routing documentation near the stage
descriptions and validate() comments to reflect lookup classification rather
than removed anchors/closed_under matching, state that lookup returns at most
one family (or none), and describe analyzer attachment as occurring during
planning after the graph is frozen. Ensure the comments no longer claim multiple
family claims compete or that validate() resolves analyzers.

In `@python/cudnn/sdpa/graph_analyzer.py`:
- Around line 545-556: Update the stale analyzer lifecycle documentation: in
python/cudnn/sdpa/graph_analyzer.py lines 545-556, describe planning followed by
_freeze() and _attach_facts() instead of validate(); make the same
planning-once-per-frozen-graph correction in lines 10-13. In
python/cudnn/sdpa/fwd/engines.py lines 482-483 and
python/cudnn/sdpa/bwd/engines.py lines 225-226, replace references to the record
validate() attached with planning-attached wording; no behavioral changes are
needed.
- Around line 46-51: Update to_torch_dtype to validate that dt exists in
_TORCH_FROM_CUDNN before indexing it, and raise the project’s typed
decline/unsupported error with the unmapped dtype included when it does not.
Preserve the existing torch dtype conversion for mapped values so
tensor_desc_from_ir can decline unsupported O or Stats dtypes instead of
propagating KeyError.

In `@test/python/test_engine_router.py`:
- Around line 799-943: Add the pytest.mark.L0 decorator to each of the seven new
test functions in this diff: test_classification_is_a_partition,
test_a_graph_spanning_two_families_belongs_to_neither,
test_family_id_blocks_are_disjoint, test_declared_analyzers_are_importable,
test_planning_attaches_facts_without_anyone_asking,
test_ranking_and_engine_read_the_same_record, and
test_facts_are_recomputed_when_the_graph_grows.

---

Nitpick comments:
In `@python/cudnn/_pygraph.py`:
- Around line 1013-1037: The _attach_facts docstring describes multi-family
discovery and attachment, but manifest.family_for() returns only one family.
Update the docstring to use singular wording throughout, stating that the graph
finds its family and attaches at most one family-scoped record, while preserving
the existing no-family behavior.
- Around line 1002-1011: Extract the duplicated backend-decline state reset into
a `_record_backend_decline(exc)` helper, including logging the exception,
storing `_backend_declined`, clearing `_lowered_graph` and `_cpp_tensors`,
resetting `_cpp_bog_done` and `_cpp_plans_created`, and emptying
`_backend_entries`. Replace the rollback blocks in the `_lower_backend_graph()`
exception path and `_finalize_backend_layout` with calls to this helper so both
paths remain synchronized.

In `@test/python/test_engine_router.py`:
- Around line 877-878: Define a module-level probe factory in
test_engine_router.py that returns an empty list, then replace the invalid
"unused_factory" value with that factory’s name in both EngineFamily
construction sites for IDs _OOT + 900 and _OOT + 920/921. Keep the existing
analyzer configuration unchanged.

🪄 Autofix

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: b929af2b-03e8-4345-b5a1-b2bf04d718a7

📥 Commits

Reviewing files that changed from the base of the PR and between 67c7634 and b5160e8.

📒 Files selected for processing (20)

• python/cudnn/_pygraph.py
• python/cudnn/engines/__init__.py
• python/cudnn/engines/engine_ids.py
• python/cudnn/engines/heuristics.py
• python/cudnn/engines/manifest.py
• python/cudnn/gemm/frost/graph_analyzer.py
• python/cudnn/linear_attention/__init__.py
• python/cudnn/linear_attention/cutile/gdn_engine.py
• python/cudnn/linear_attention/cutile/kda_engine.py
• python/cudnn/linear_attention/frost/gdn2_engine.py
• python/cudnn/linear_attention/frost/gdn_engine.py
• python/cudnn/linear_attention/frost/kda_engine.py
• python/cudnn/sdpa/__init__.py
• python/cudnn/sdpa/bwd/__init__.py
• python/cudnn/sdpa/bwd/engines.py
• python/cudnn/sdpa/fwd/__init__.py
• python/cudnn/sdpa/fwd/engines.py
• python/cudnn/sdpa/graph_analyzer.py
• test/python/sdpa/frost/test_sdpa_graph_analyzer.py
• test/python/test_engine_router.py

[coderabbitai]

Actionable comments posted: 3

🤖 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 `@test/python/test_engine_router.py`:
- Around line 909-911: Apply Black formatting to both EngineFamily fixture
constructors in test/python/test_engine_router.py at lines 909-911 and 954-956,
wrapping their arguments across multiple lines while preserving the existing
values and behavior.

In `@test/python/test_import_boundaries.py`:
- Line 36: Rename the ambiguous loop variable in the stdout line-selection
expression to a descriptive name such as line, updating its references
consistently so Ruff E741 passes.
- Around line 33-35: Update the probe execution handling around subprocess.run
so a non-zero return code fails the test instead of calling pytest.skip. Keep
skipping only for an unavailable probe interpreter if that case is explicitly
distinguishable, and ensure failures from the probe code, including import
errors, are surfaced with stderr details.

🪄 Autofix

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: ba1c1344-c04f-4a81-8754-0dafdf00e8ab

📥 Commits

Reviewing files that changed from the base of the PR and between bc68048 and a277c14.

📒 Files selected for processing (22)

• python/cudnn/_pygraph.py
• python/cudnn/engines/manifest.py
• python/cudnn/gemm/frost/engine.py
• python/cudnn/linear_attention/__init__.py
• python/cudnn/linear_attention/cutile/gdn_engine.py
• python/cudnn/linear_attention/cutile/kda_engine.py
• python/cudnn/linear_attention/frost/gdn2_engine.py
• python/cudnn/linear_attention/frost/gdn_engine.py
• python/cudnn/linear_attention/frost/kda_engine.py
• python/cudnn/sdpa/bwd/engine.py
• python/cudnn/sdpa/bwd/engines.py
• python/cudnn/sdpa/fwd/engine.py
• python/cudnn/sdpa/fwd/engines.py
• test/python/gemm/frost/test_frontend_integration.py
• test/python/linear_attention/frost/test_gdn2_bprop_kernel.py
• test/python/linear_attention/frost/test_gdn2_prefill_kernel.py
• test/python/linear_attention/frost/test_gdn_bprop_kernel.py
• test/python/linear_attention/frost/test_gdn_prefill_kernel.py
• test/python/linear_attention/frost/test_kda_bprop_kernel.py
• test/python/linear_attention/frost/test_kda_prefill_kernel.py
• test/python/test_engine_router.py
• test/python/test_import_boundaries.py

💤 Files with no reviewable changes (12)

• test/python/linear_attention/frost/test_gdn2_bprop_kernel.py
• python/cudnn/linear_attention/frost/gdn_engine.py
• python/cudnn/linear_attention/frost/kda_engine.py
• test/python/linear_attention/frost/test_gdn_prefill_kernel.py
• python/cudnn/linear_attention/frost/gdn2_engine.py
• python/cudnn/linear_attention/cutile/kda_engine.py
• python/cudnn/linear_attention/cutile/gdn_engine.py
• test/python/linear_attention/frost/test_gdn2_prefill_kernel.py
• test/python/linear_attention/frost/test_kda_prefill_kernel.py
• python/cudnn/_pygraph.py
• test/python/linear_attention/frost/test_kda_bprop_kernel.py
• test/python/linear_attention/frost/test_gdn_bprop_kernel.py

🚧 Files skipped from review as they are similar to previous changes (4)

• python/cudnn/linear_attention/init.py
• python/cudnn/sdpa/bwd/engines.py
• python/cudnn/sdpa/fwd/engines.py
• python/cudnn/engines/manifest.py

[coderabbitai]

📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win

Format the EngineFamily fixture constructors with Black.

These constructors exceed the configured 160-character line length.

• test/python/test_engine_router.py#L909-L911: Wrap the EngineFamily arguments with Black formatting.
• test/python/test_engine_router.py#L954-L956: Apply the same formatting to the duplicate fixture constructor.

📍 Affects 1 file

• test/python/test_engine_router.py#L909-L911 (this comment)
• test/python/test_engine_router.py#L954-L956

🤖 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_engine_router.py` around lines 909 - 911, Apply Black
formatting to both EngineFamily fixture constructors in
test/python/test_engine_router.py at lines 909-911 and 954-956, wrapping their
arguments across multiple lines while preserving the existing values and
behavior.

Source: Coding guidelines

[YangXu1990uiuc]

Thanks — verified each against HEAD rather than the commit it was filed on, since three were already addressed by later commits in the series.

Fixed (d6e2b01f)

• to_torch_dtype unguarded on an unmapped type. Correct and the most useful of the six. Only Q's dtype is capability-checked, and tensor_desc_from_ir() runs on every bound tensor, so O / Stats / a side output can reach it with a type outside the map — KeyError deep in lowering instead of a decline. It raises NotImplementedError now, which the router treats as "this engine does not serve this graph".
• Docstrings still naming validate() as where facts are attached. Two were left (manifest.py, graph_analyzer.py module docstring). Planning attaches them, after _finalize_backend_layout() and _freeze().

Already fixed before the review landed — these were filed against b5160e860 / bc68048dd:

• The _facts comment describing a (node count, facts) tuple: it is # analyzer callable -> its record; see _facts_for() as of a277c141, which also changed the key from a module.qualname string to the callable itself.
• The manifest module docstring describing anchors / closed_under matching and competing family claims: rewritten in b5160e86. The one surviving closed_under mention is a deliberate note about why it was deleted.
• Test-level markers on the new tests: both files carry a module-level pytestmark = pytest.mark.L0, which the new tests inherit.

Not doing: the EngineRow compatibility export.

EngineRow is not a compatibility surface. It arrived with #476, which merged into develop on 2026-08-04 — two days before this PR — and has never been in a release, so nothing downstream can be importing it. It also has no use outside the library: it describes an in-tree engine family for the manifest to classify graphs with, while the out-of-tree surface is BaseEngine / Router / register_backend(), all unchanged.

More to the point, a compatibility adapter could not preserve the old behaviour even if we wanted it to. EngineRow carried anchors and closed_under, and both are gone by design: classification is a lookup table now (a graph belongs to one family or none, rather than N families each declaring a claim), and closed_under was a coarser duplicate of check_support() that promised RESHAPE support nothing implemented. A shim returning an object with those fields would be a fresh place to lie.

The PR description records this under API impact.

[coderabbitai]

Actionable comments posted: 1

🤖 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/sdpa/graph_analyzer.py`:
- Around line 11-14: Update the module docstring wording near the SDPA family
names to use the grammatically complete phrase “callable for the SDPA family
names,” without changing the surrounding documentation.

🪄 Autofix

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: 2d9b674c-7180-4b3b-9460-f4a73a995d60

📥 Commits

Reviewing files that changed from the base of the PR and between a277c14 and d6e2b01.

📒 Files selected for processing (2)

• python/cudnn/engines/manifest.py
• python/cudnn/sdpa/graph_analyzer.py

🚧 Files skipped from review as they are similar to previous changes (1)

• python/cudnn/engines/manifest.py

[coderabbitai]

📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win

Fix the module docstring grammar.

The phrase callable the SDPA family names is incomplete. Rewrite it so the relationship is explicit, for example by changing it to callable for the SDPA family names.

Proposed fix

-callable the SDPA family names in ``engines/manifest.py``; PLANNING runs it once
+callable for the SDPA family names in ``engines/manifest.py``; PLANNING runs it once

📝 Committable suggestion

‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.

Suggested change

	callable the SDPA family names in ``engines/manifest.py``; PLANNING runs it once
	per graph -- after the backend's layout inference has landed and the graph is
	frozen -- and attaches the record, so the ranking and the engine share it rather
	than each parsing the graph.
	callable for the SDPA family names in ``engines/manifest.py``; PLANNING runs it once
	per graph -- after the backend's layout inference has landed and the graph is
	frozen -- and attaches the record, so the ranking and the engine share it rather
	than each parsing the graph.

🤖 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/sdpa/graph_analyzer.py` around lines 11 - 14, Update the module
docstring wording near the SDPA family names to use the grammatically complete
phrase “callable for the SDPA family names,” without changing the surrounding
documentation.

[YangXu1990uiuc]

Latest round — two fixed, two not applicable.

Fixed

• pytest.skip hid a failing probe. Good catch: a probe that cannot import cudnn is exactly the regression that file exists to catch, and skipping reported it as "not checked". pytest.fail now.
• Ruff E741 l → ln.

Not applicable

• Black on the EngineFamily fixtures. Neither construct exceeds 160 chars, and black --line-length 160 --check reports both files unchanged. pre-commit runs black on every commit in this branch.
• "callable the SDPA family names" grammar. names is the verb, not a noun: the callable [that] the SDPA family names in engines/manifest.py — a reduced relative clause. The proposed "callable for the SDPA family names" reparses names as a noun and changes the meaning.

Also worth flagging, since it came out of a question about this PR rather than from review: deferring the DSL import past check_support had turned "cutedsl extra not installed" from the family is absent into an ImportError out of build_plan() — which is in neither the engine's except clause nor decline_types(), so it propagated past the backend sitting in the same ranked list. The engine now declines at check time using probes that never execute the module, and ImportError joins the decline types as a second line. Regression test added.

[vedaanta]

do we need this?

[YangXu1990uiuc]

No — you're right, and it's reverted in the next push.

Pinning the extra to >=4.7.0 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. The extra stays >=4.5.0.

The 4.7.0 floor is real — it's what the CuTe primitives these engines lower through need — but it belongs at runtime, not in the dependency:

• engine: check_support() declines when the installed DSL is older, so the graph goes to the backend and only those engines are lost.
• tests: the same predicate, so an older DSL makes the suite skip rather than fail.

Both read CUTEDSL_MIN_VERSION from cudnn/frost/buffers.py, so there's one number rather than two that can drift.

The probe doesn't cost anything either — importlib.util.find_spec (7 ms) and importlib.metadata.version (5 ms), neither of which executes the module, against ~4.3 s and 1410 modules for a real import cutlass.

[YangXu1990uiuc]

/bot run

[Anerudhan]

@cudnn-ci-bot run frost

[cudnn-ci-bot]

🚀 Running mirror pipeline

Branch: cudnn-gh/pr-502-9c4fa64
Pipeline: 61528884
Targets: frost

[Anerudhan]

Few things:

1. Can you check-in a design doc so it is easier for agents in future..
2. Can you update your agents to create labels for your PR..
3. Use the PR template for license.

Aside on the code:

1. I am still vague about how the stream/current_device works. Is it a fallback if user does not provide the device?.

[YangXu1990uiuc]

All four — thanks.

Design doc. There is one, docs/python_graph_and_execution_backends.md, checked in since #336 — it just described the shape this PR replaces, so following it you'd have gone looking for anchors and closed_under that no longer exist. Updated here (cee6a0a3) rather than starting a second one, since two design docs is exactly the duplicate-that-can-lie problem the rest of this PR is deleting.

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 resolve; and which imports each dispatch stage may pay for. Existing sections corrected where this PR moved them.

Labels. Added: cat-cleanup, mod-frontend, mod-frost, orig-nv-eng (and on #506: cat-bug, mod-frost, mod-cutedsl, orig-nv-eng). I'll make this part of opening a PR rather than a follow-up.

PR template. My fault — I wrote the body from scratch and dropped the checklist. Restored, license included.

Stream / current_device. Three separate things, and the short answer to your question is: it's a fallback for having no handle, not for a handle we failed to read.

• handle — execute(..., handle=h) if given, else the graph's own. The handle is what carries the stream.
• stream — _resolve_stream(handle) is cudnn.get_stream(handle). No handle → None, and kernel wrappers then use the default stream. But a failed query on a supplied handle raises rather than quietly falling back: running on the wrong stream is a correctness bug, not a degradation, so it must not be silently absorbed.
• device — this PR changes it. The analyzer used to read compute capability and SM count from torch.cuda.current_device(), i.e. ambient state that a graph carries around and that goes stale if you switch device or pass a handle for another one. It now reads cudnn.create_device_properties() — the backend's own serialisable descriptor, the same object the C++ deviceless-AoT path uses. Engines re-check the arch in check_support() regardless.

That last one has a second payoff worth flagging: facts computed from a serialisable descriptor need no live device, which is the precondition for a python engine to participate in deviceless AoT at all. It cannot today.

Relatedly, #506 (separate PR, off develop) fixes a stream bug this investigation turned up: buffers.probe() called __dlpack__() with the default stream argument, so torch did a record_stream that is illegal inside a CUDA graph capture. 286 linear-attention tests, red on develop today.