Four graph-API bugs, each with the regression test it needed

python/cudnn/__init__.py

@@ -381,23 +381,17 @@ def __getattr__(name: str) -> Any:
         globals()["Graph"] = _wrapper.Graph
         return globals()[name]
 
-    if name == "ops":
-        # Use importlib rather than "from . import ops" to avoid infinite
+    if name in ("ops", "experimental"):
+        # Use importlib rather than "from . import <name>" to avoid infinite
         # recursion. The cycle:
-        #   1. cudnn.ops accessed → __getattr__("ops") fires
-        #   2. "from . import ops" → _handle_fromlist(cudnn, ["ops"], ...)
-        #   3. _handle_fromlist calls hasattr(cudnn, "ops")
-        #   4. "ops" not in __dict__ yet → __getattr__("ops") again → goto 1
+        #   1. cudnn.<name> accessed → __getattr__("<name>") fires
+        #   2. "from . import <name>" → _handle_fromlist(cudnn, ["<name>"], ...)
+        #   3. _handle_fromlist calls hasattr(cudnn, "<name>")
+        #   4. not in __dict__ yet → __getattr__("<name>") again → goto 1
         # importlib.import_module bypasses _handle_fromlist entirely.
-        _ops = importlib.import_module(".ops", __name__)
-        globals()["ops"] = _ops
-        return _ops
-
-    if name == "experimental":
-        from . import experimental as _experimental
-
-        globals()["experimental"] = _experimental
-        return _experimental
+        module = importlib.import_module(f".{name}", __name__)
+        globals()[name] = module
+        return module
 
     if name in _LAZY_OPTIONAL_IMPORTS:
         return _load_optional_symbol(name)

python/cudnn/frost/device.py

@@ -49,7 +49,7 @@ def _device_handle(device: int):
     """``CUdevice`` for an ordinal. Needs only cuInit — creates no context."""
     drv = _driver()
     if drv is None:
-        raise RuntimeError("cudnn.frost: no CUDA device visible")
+        raise NotImplementedError("cudnn.frost: no CUDA device visible")
     count = int(_ck(*drv.cuDeviceGetCount()))
     if not 0 <= device < count:
         raise ValueError(f"cudnn.frost: cuda:{device} does not exist ({count} device(s) visible)")
@@ -65,7 +65,7 @@ def current_device() -> int:
     slot, so that is the second rung."""
     drv = _driver()
     if drv is None:
-        raise RuntimeError("cudnn.frost: no CUDA device visible")
+        raise NotImplementedError("cudnn.frost: no CUDA device visible")
     if int(_ck(*drv.cuCtxGetCurrent())) != 0:
         return int(_ck(*drv.cuCtxGetDevice()))
     import cuda.bindings.runtime as rt
@@ -174,7 +174,7 @@ def __init__(self, device: int):
     def __enter__(self):
         self._drv = _driver()
         if self._drv is None:
-            raise RuntimeError("cudnn.frost: no CUDA device visible")
+            raise NotImplementedError("cudnn.frost: no CUDA device visible")
         self._handle = _device_handle(self._device)
         self._previous = _ck(*self._drv.cuCtxGetCurrent())
         _ck(*self._drv.cuCtxSetCurrent(_ck(*self._drv.cuDevicePrimaryCtxRetain(self._handle))))

python/cudnn/frost/workspace.py

@@ -4,7 +4,7 @@
 """Per-execute scratch carved out of the CALLER's workspace — shared by every
 FROST engine.
 
-The workspace contract (see ``cudnn/frost/dispatch.py``) is that an executor
+The workspace contract (see ``cudnn/engines/base.py``) is that an executor
 never allocates: it reports the scratch it needs as ``workspace_bytes`` and
 carves that scratch out of the buffer ``execute()`` hands it, so the pointers
 are caller-owned and stable across executes — which is what makes a plan safe

python/cudnn/gemm/frost/compiler.py

@@ -1794,9 +1794,10 @@ class CompiledFusedGemm:
     # clamped); drives the runtime output/aux alignment requirements. None →
     # fall back to the chain-derived width.
     vec_bytes_epi: "int | None" = None
-    # Opt in to stream-aware dispatch: frost/dispatch.py resolves the stream
-    # from the execute-time cuDNN handle and forwards it as `stream=`. Engines
-    # that do not carry the param stay on the default stream (see dispatch).
+    # Opt in to stream-aware dispatch: the engine resolves the stream from the
+    # execute-time cuDNN handle into ExecutionContext.stream (cudnn/engines/base.py)
+    # and forwards it as `stream=`. Engines that do not carry the param stay on
+    # the default stream.
     accepts_stream: ClassVar[bool] = True
 
     def __call__(self, variant_pack, stream=None):

python/cudnn/gemm/frost/tile_config.py

@@ -25,10 +25,10 @@ def _sm_smem_budget_bytes_of(device: int) -> int:
     from cudnn.frost.device import device_name, is_available, shared_memory_per_block_optin
 
     if not is_available():
-        raise RuntimeError("cannot size the SMEM pipeline: no CUDA device is visible to query MaxSharedMemoryPerBlockOptin")
+        raise NotImplementedError("cannot size the SMEM pipeline: no CUDA device is visible to query MaxSharedMemoryPerBlockOptin")
     optin = shared_memory_per_block_optin(device)
     if not optin:
-        raise RuntimeError(f"the driver did not report MaxSharedMemoryPerBlockOptin for device {device_name(device)!r}; cannot size the SMEM pipeline")
+        raise NotImplementedError(f"the driver did not report MaxSharedMemoryPerBlockOptin for device {device_name(device)!r}; cannot size the SMEM pipeline")
     return int(optin)
 
 
@@ -57,10 +57,10 @@ def _l2_swizzle_budget_bytes_of(device: int) -> int:
     from cudnn.frost.device import device_name, is_available, l2_cache_bytes
 
     if not is_available():
-        raise RuntimeError("cannot size the L2 tile-rasterization budget: no CUDA device is visible to query L2CacheSize")
+        raise NotImplementedError("cannot size the L2 tile-rasterization budget: no CUDA device is visible to query L2CacheSize")
     l2 = l2_cache_bytes(device)
     if not l2:
-        raise RuntimeError(f"the driver did not report L2CacheSize for device {device_name(device)!r}; cannot size the L2 tile-rasterization budget")
+        raise NotImplementedError(f"the driver did not report L2CacheSize for device {device_name(device)!r}; cannot size the L2 tile-rasterization budget")
     return int(l2) // _L2_RETENTION_DIVISOR
 
 

python/cudnn/linear_attention/cutile/gdn_engine.py

@@ -125,7 +125,7 @@ def check_support(self, graph: "pygraph") -> None:
             if int(err) != 0:
                 raise NotImplementedError(f"GdnCuTileEngine: cudaRuntimeGetVersion failed ({err})")
         except ImportError as e:
-            raise NotImplementedError(f"GdnCuTileEngine requires cuda.bindings: {e}")
+            raise NotImplementedError(f"GdnCuTileEngine: 'from cuda.bindings import runtime' failed ({e})")
         if _cudart_version < 13030:
             raise NotImplementedError(f"GdnCuTileEngine requires CUDA 13.3+ (found {_cudart_version})")
         try:

python/cudnn/linear_attention/cutile/kda_engine.py

@@ -137,7 +137,7 @@ def check_support(self, graph: "pygraph") -> None:
             if int(err) != 0:
                 raise NotImplementedError(f"KdaCuTileEngine: cudaRuntimeGetVersion failed ({err})")
         except ImportError as e:
-            raise NotImplementedError(f"KdaCuTileEngine requires cuda.bindings: {e}")
+            raise NotImplementedError(f"KdaCuTileEngine: 'from cuda.bindings import runtime' failed ({e})")
         if _cudart_version < 13030:
             raise NotImplementedError(f"KdaCuTileEngine requires CUDA 13.3+ (found {_cudart_version})")
         try:

python/cudnn/linear_attention/frost/gdn2_engine.py

@@ -95,7 +95,7 @@ def check_support(self, graph) -> None:
         try:
             import cutlass.experimental.primitives  # noqa: F401 — availability probe: ImportError = decline
         except ImportError as exc:
-            raise NotImplementedError(f"Gdn2FrostEngine requires the Cutlass DSL with cutlass.experimental.primitives: {exc}") from exc
+            raise NotImplementedError(f"Gdn2FrostEngine: 'import cutlass.experimental.primitives' failed ({exc})") from exc
         for port in ("q", "k", "v", "g", "beta", "w", "cu_seqlens"):
             if port not in node.inputs:
                 raise NotImplementedError(f"Gdn2FrostEngine: GDN2 node '{node.name}' is missing input '{port}'")

python/cudnn/linear_attention/frost/gdn_engine.py

@@ -117,7 +117,7 @@ def check_support(self, graph) -> None:
         try:
             import cutlass.experimental.primitives  # noqa: F401 — availability probe: ImportError = decline
         except ImportError as exc:
-            raise NotImplementedError(f"GdnFrostEngine requires the Cutlass DSL with cutlass.experimental.primitives: {exc}") from exc
+            raise NotImplementedError(f"GdnFrostEngine: 'import cutlass.experimental.primitives' failed ({exc})") from exc
         if node.params.get("use_qk_l2norm", False):
             raise NotImplementedError("GdnFrostEngine: use_qk_l2norm is not supported (the kernel takes q/k as given)")
         ports = ("q", "k", "v", "g", "beta", "cu_seqlens")

python/cudnn/linear_attention/frost/kda_engine.py

@@ -98,7 +98,7 @@ def check_support(self, graph) -> None:
         try:
             import cutlass.experimental.primitives  # noqa: F401 — availability probe: ImportError = decline
         except ImportError as exc:
-            raise NotImplementedError(f"KdaFrostEngine requires the Cutlass DSL with cutlass.experimental.primitives: {exc}") from exc
+            raise NotImplementedError(f"KdaFrostEngine: 'import cutlass.experimental.primitives' failed ({exc})") from exc
         for port in ("q", "k", "v", "g", "beta", "cu_seqlens"):
             if port not in node.inputs:
                 raise NotImplementedError(f"KdaFrostEngine: KDA node '{node.name}' is missing input '{port}'")

python/cudnn/wrapper.py

@@ -220,6 +220,11 @@ class Graph:
     """
 
     __handle: Optional[CudnnHandle] = None  # holding the cudnn handle pointer
+    # None: allocate one on __exit__. False: the caller owns it (workspace_alloc=False),
+    # and execute() demands a workspace= kwarg. A class attribute, not a hasattr()
+    # probe: "__workspace" inside hasattr is a plain string, so it is NOT name-mangled
+    # and never matched the _Graph__workspace the assignments below produce.
+    __workspace: Any = None
 
     def __init__(
         self,
@@ -305,7 +310,7 @@ def __exit__(self, exc_type, exc_value, tb):
         self.__graph.check_support()
         self.__graph.build_plans()
         # Set up workspace if not forbidden by user, then set up I/O tensor orders
-        if not hasattr(self, "__workspace"):
+        if self.__workspace is None:
             self.__workspace = torch.empty(
                 self.__graph.get_workspace_size(),
                 device="cuda",

test/python/test_decline_types.py

@@ -0,0 +1,65 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
+"""An engine that cannot serve a graph must say so with a decline type.
+
+``build_plans()`` walks the ranked plan list and skips an entry that raises one
+of ``engines.base.decline_types()``, moving on to the next plan and ultimately
+to the cuDNN backend. Anything else propagates and aborts the walk, so a graph
+the backend could have served fails outright.
+
+"This machine has no CUDA device" and "the driver did not report the property I
+need to size a pipeline" are declines: the engine cannot serve the graph, but
+another entry in the list can. They were raising RuntimeError, which is not a
+decline type and is not caught by the engines' own ``build_plan`` handlers
+either, so a probe failure took down the whole walk instead of falling back.
+"""
+
+import pytest
+
+import cudnn
+from cudnn.engines.base import decline_types
+
+
+@pytest.mark.L0
+def test_device_probes_decline_when_no_driver(monkeypatch):
+    from cudnn.frost import device
+
+    monkeypatch.setattr(device, "_driver", lambda: None)
+
+    with pytest.raises(decline_types()):
+        device.current_device()
+    with pytest.raises(decline_types()):
+        device._device_handle(0)
+    with pytest.raises(decline_types()):
+        with device.device_context(0):
+            pass
+
+
+@pytest.mark.L0
+@pytest.mark.parametrize("probe", ["_sm_smem_budget_bytes_of", "_l2_swizzle_budget_bytes_of"])
+def test_tile_config_probes_decline_when_unavailable(monkeypatch, probe):
+    from cudnn.frost import device as frost_device
+    from cudnn.gemm.frost import tile_config
+
+    fn = getattr(tile_config, probe)
+    # Both probes are @lru_cache'd, so an earlier test that already queried this
+    # device would serve a cached answer and never reach the raise.
+    fn.cache_clear()
+    monkeypatch.setattr(frost_device, "is_available", lambda: False)
+    try:
+        with pytest.raises(decline_types()):
+            fn(0)
+    finally:
+        fn.cache_clear()
+
+
+@pytest.mark.L0
+def test_decline_types_are_what_build_plans_skips():
+    """The tuple is the contract; keep it and the walk in agreement."""
+    assert NotImplementedError in decline_types()
+    assert cudnn.cudnnGraphNotSupportedError in decline_types()
+    assert ImportError in decline_types()
+    # RuntimeError must NOT be a decline: it is how an engine reports a bug,
+    # and swallowing it would hide real failures behind a silent fallback.
+    assert RuntimeError not in decline_types()

test/python/test_import_boundaries.py

@@ -97,3 +97,22 @@ def test_support_check_pulls_no_framework(module):
     process that merely asks whether an engine applies.
     """
     _assert_absent(_imported_by(f"import cudnn\nimport {module}"), module)
+
+
+@pytest.mark.parametrize("name", ["ops", "experimental", "wrapper", "Graph"])
+def test_lazy_top_level_attribute_resolves(name):
+    """Every name cudnn.__getattr__ special-cases must actually resolve.
+
+    Regression: `experimental` was fetched with `from . import experimental`,
+    the one form the comment six lines above it documents as recursive —
+    _handle_fromlist calls hasattr(cudnn, "experimental"), which re-enters
+    __getattr__ because the name is not in __dict__ yet. `cudnn.experimental`
+    raised RecursionError, so the documented experimental namespace was
+    unreachable by attribute access.
+    """
+    out = subprocess.run(
+        [sys.executable, "-c", f"import cudnn; x = cudnn.{name}; print(type(x).__name__)"],
+        capture_output=True,
+        text=True,
+    )
+    assert out.returncode == 0, f"cudnn.{name} failed:\n{out.stderr[-1500:]}"

test/python/test_wrapper_graph.py

@@ -0,0 +1,57 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+
+"""Tests for the fluent ``cudnn.Graph`` wrapper (python/cudnn/wrapper.py)."""
+
+import pytest
+import torch
+
+import cudnn
+
+
+def _matmul_graph(**kwargs):
+    """A 64x64 half matmul through the fluent wrapper."""
+    with cudnn.Graph(
+        handle="auto",
+        io_data_type=cudnn.data_type.HALF,
+        compute_data_type=cudnn.data_type.FLOAT,
+        inputs=["X", "W"],
+        outputs=["Y"],
+        **kwargs,
+    ) as graph:
+        X = graph.tensor(name="X", dim=[1, 64, 64], stride=[64 * 64, 64, 1])
+        W = graph.tensor(name="W", dim=[1, 64, 64], stride=[64 * 64, 64, 1])
+        Y = graph.matmul(name="mm", A=X, B=W)
+        Y.set_output(True).set_name("Y")
+    return graph
+
+
+@pytest.mark.L0
+def test_workspace_alloc_default_allocates():
+    """The default path allocates a workspace the caller never has to think about."""
+    graph = _matmul_graph()
+    assert torch.is_tensor(graph._Graph__workspace)
+
+
+@pytest.mark.L0
+def test_workspace_alloc_false_is_honored():
+    """``workspace_alloc=False`` means the CALLER owns the workspace.
+
+    Regression: the sentinel is written as ``self.__workspace`` (mangled to
+    ``_Graph__workspace``) but was read back with ``hasattr(self, "__workspace")``
+    — a plain string, which is NOT name-mangled. That probe was therefore always
+    False, the sentinel was overwritten with a fresh allocation on every
+    ``__exit__``, and the "Need to specify workspace" guard below was unreachable.
+    """
+    graph = _matmul_graph(workspace_alloc=False)
+    assert graph._Graph__workspace is False
+
+    x = torch.randn(1, 64, 64, dtype=torch.half, device="cuda")
+    w = torch.randn(1, 64, 64, dtype=torch.half, device="cuda")
+
+    with pytest.raises(RuntimeError, match="Need to specify workspace"):
+        graph(x, w)
+
+    workspace = torch.empty(max(graph.get_workspace_size(), 1), dtype=torch.uint8, device="cuda")
+    out = graph(x, w, workspace=workspace)
+    torch.testing.assert_close(out.float(), (x @ w).float(), atol=1e-2, rtol=1e-2)