Publish the graph API as cuda.core.graph (#1858)

* Reorganize graph test files for clarity

Rename test files to reflect what they actually test:
- test_basic -> test_graph_builder (stream capture tests)
- test_conditional -> test_graph_builder_conditional
- test_advanced -> test_graph_update (moved child_graph and
  stream_lifetime tests into test_graph_builder)
- test_capture_alloc -> test_graph_memory_resource
- test_explicit* -> test_graphdef*

Made-with: Cursor

* Enhance Graph.update() and add whole-graph update tests

- Extend Graph.update() to accept both GraphBuilder and GraphDef sources
- Surface CUgraphExecUpdateResultInfo details on update failure instead
  of a generic CUDA_ERROR_GRAPH_EXEC_UPDATE_FAILURE message
- Release the GIL during cuGraphExecUpdate via nogil block
- Add parametrized happy-path test covering both GraphBuilder and GraphDef
- Add error-case tests: unfinished builder, topology mismatch, wrong type

Made-with: Cursor

* Add AdjacencySet proxy for pred/succ and GraphNode.remove()

Replace cached tuple-based pred/succ with mutable AdjacencySet backed
by direct CUDA driver calls. Add GraphNode.remove() wrapping
cuGraphDestroyNode.

Made-with: Cursor

* Add edge mutation support and MutableSet interface for GraphNode adjacencies

Enable adding/removing edges between graph nodes via AdjacencySet (a
MutableSet proxy on GraphNode.pred/succ), node removal via discard(),
and property setters for bulk edge replacement. Includes comprehensive
mutation and interface tests.

Closes part of #1330 (step 2: edge mutation on GraphDef).

Made-with: Cursor

* Use requires_module mark for numpy version checks in mutation tests

Replace inline skipif version check with requires_module(np, "2.1")
from the shared test helpers, consistent with other test files.

Made-with: Cursor

* Fix empty-graph return type: return set() instead of () for nodes/edges

Made-with: Cursor

* Rename AdjacencySet to AdjacencySetProxy, add bulk ops and safety guards

Rename class and file to AdjacencySetProxy to clarify write-through
semantics. Add bulk-efficient clear(), __isub__(), __ior__() overrides
and remove_edges() on the Cython core. Guard GraphNode.discard() against
double-destroy via membership check. Filter duplicates in update(). Add
error-path tests for wrong types, cross-graph edges, and self-edges.

Made-with: Cursor

* Add destroy() method with handle invalidation, remove GRAPH_NODE_SENTINEL

Replace discard() with destroy() which calls cuGraphDestroyNode and then
zeroes the CUgraphNode resource in the handle box via
invalidate_graph_node_handle. This prevents stale memory access on
destroyed nodes. Properties (type, pred, succ, handle) degrade gracefully
to None/empty for destroyed nodes.

Remove the GRAPH_NODE_SENTINEL (0x1) approach in favor of using NULL for
both sentinels and destroyed nodes, which is simpler and avoids the risk
of passing 0x1 to driver APIs that treat it as a valid pointer.

Made-with: Cursor

* Add GraphNode identity cache for stable Python object round-trips

Nodes retrieved via GraphDef.nodes(), edges(), or pred/succ traversal
now return the same Python object that was originally created, enabling
identity checks with `is`. A C++ HandleRegistry deduplicates
CUgraphNode handles, and a Cython WeakValueDictionary caches the
Python wrapper objects.

Made-with: Cursor

* Purge node cache on destroy to prevent stale identity lookups

Made-with: Cursor

* Make graph API public: rename _graph to cuda.core.graph

Move the graph package from cuda.core._graph to cuda.core.graph and
flatten the _graph_def subdirectory. GraphDef, Condition, GraphAllocOptions,
GraphNode, and all node subclasses are now importable from cuda.core.graph.
GraphDef, Condition, and GraphAllocOptions are also re-exported from
cuda.core directly.

Break the circular import (_stream → graph → _graph_node →
_kernel_arg_handler → _buffer → _device → _stream) by making _device.pyx
and _stream.pyx use local imports for GraphBuilder.

Made-with: Cursor

* Add 0.7.x release notes for GraphDef and cuda.core.graph module

Made-with: Cursor

* Skip NULL nodes in graph_node_registry to fix sentinel identity collision

Sentinel (entry) nodes use NULL as their CUgraphNode, so caching them
under a NULL key caused all sentinels across different graphs to share
the same handle. This made nodes built from the wrong graph's entry
point, causing CUDA_ERROR_INVALID_VALUE for conditional nodes and hash
collisions in equality tests.

Made-with: Cursor

* Unregister destroyed nodes from C++ graph_node_registry

When a node is destroyed, the driver may reuse its CUgraphNode pointer
for a new node. Without unregistering the old entry, the registry
returns a stale handle pointing to the wrong node type and graph.

Made-with: Cursor

* Add dedicated test for node identity preservation through round-trips

Made-with: Cursor

* Add API docs for cuda.core.graph and fix stale docstring references

The _graph -> graph rename left behind broken Sphinx cross-references
and the new graph types were missing from the API reference.

Made-with: Cursor

* Add handle= to all GraphNode subclass __repr__ for debugging

Every subclass repr now starts with handle=0x... (the CUgraphNode
pointer) followed by type-specific identity/parameter data. Dynamic
queries (pred counts, subnode counts) are removed in favor of
deterministic, cheap fields. This makes set comparison failures in
test output readable when debugging graph mutation tests.

Made-with: Cursor

* Rename _node_cache/_cached to _node_registry/_registered

Aligns Python-side terminology with the C++ graph_node_registry.

Made-with: Cursor

* Fix unregister_handle and rename invalidate_graph_node_handle

unregister_handle: remove the expired() guard that prevented erasure
when the shared_ptr was still alive. This caused stale registry
entries after destroy(), leading to CUDA_ERROR_INVALID_VALUE when
the driver reused CUgraphNode pointer values.

Rename invalidate_graph_node_handle -> invalidate_graph_node for
consistency with the rest of the graph node API.

Made-with: Cursor

* Add cheap containment test and early type check for AdjacencySetProxy

Add _AdjacencySetCore.contains() that checks membership by comparing
raw CUgraphNode handles at the C level, avoiding Python object
construction. Uses a 16-element stack buffer for a single driver call
in the common case.

Move the type check in update() inline next to the extend loop so
invalid input is rejected immediately.

Made-with: Cursor

* Add GraphDef.empty(), stack-buffer query optimization, and registry test

- Add GraphDef.empty() for creating entry-point empty nodes; replace
  all no-arg join() calls on GraphDef with empty() in tests.
- Optimize _AdjacencySetCore.query() to use a 16-element stack buffer,
  matching the contains() optimization.
- Add test_registry_cleanup exercising destroy(), graph deletion, and
  weak-reference cleanup of the node registry.

Made-with: Cursor

* Document the two-level handle and object registry design

Add REGISTRY_DESIGN.md explaining how the C++ HandleRegistry (Level 1)
and Cython _node_registry (Level 2) work together to preserve Python
object identity through driver round-trips. Add cross-references at
each registry instantiation site.

Made-with: Cursor

* Fix import formatting in test_registry_cleanup

Made-with: Cursor

* Optimize GraphDef.nodes() and edges() to try a single driver call

Pre-allocate vectors to 128 entries and pass them on the first call.
Only fall back to a second call if the graph exceeds 128 nodes/edges.

Made-with: Cursor

* Update docstrings, api.rst structure, and release notes

- Rename "Memory" section to "Memory management" in api.rst
- Add introductory paragraph under "Node types" subsection
- Update node docstrings to concise noun phrases
- Update graph API class docstrings (GraphBuilder, Graph, etc.)
- Revise release notes highlights and new features sections
- Fix test_registry_cleanup import path (cuda.core._graph -> cuda.core.graph)

Made-with: Cursor

Author: Andy-Jost

cuda_core/cuda/core/__init__.py

@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2024-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 #
 # SPDX-License-Identifier: Apache-2.0
 
@@ -31,12 +31,6 @@
 from cuda.core import system, utils
 from cuda.core._device import Device
 from cuda.core._event import Event, EventOptions
-from cuda.core._graph import (
-    Graph,
-    GraphBuilder,
-    GraphCompleteOptions,
-    GraphDebugPrintOptions,
-)
 from cuda.core._graphics import GraphicsResource
 from cuda.core._launch_config import LaunchConfig
 from cuda.core._launcher import launch
@@ -69,3 +63,12 @@
     StreamOptions,
 )
 from cuda.core._tensor_map import TensorMapDescriptor, TensorMapDescriptorOptions
+from cuda.core.graph import (
+    Condition,
+    Graph,
+    GraphAllocOptions,
+    GraphBuilder,
+    GraphCompleteOptions,
+    GraphDebugPrintOptions,
+    GraphDef,
+)

cuda_core/cuda/core/_device.pyx

@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2024-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 #
 # SPDX-License-Identifier: Apache-2.0
 
@@ -24,7 +24,6 @@ from cuda.core._resource_handles cimport (
     as_cu,
 )
 
-from cuda.core._graph import GraphBuilder
 from cuda.core._stream import IsStreamT, Stream, StreamOptions
 from cuda.core._utils.clear_error_support import assert_type
 from cuda.core._utils.cuda_utils import (
@@ -1363,15 +1362,17 @@ class Device:
         self._check_context_initialized()
         handle_return(runtime.cudaDeviceSynchronize())
 
-    def create_graph_builder(self) -> GraphBuilder:
-        """Create a new :obj:`~_graph.GraphBuilder` object.
+    def create_graph_builder(self) -> "GraphBuilder":
+        """Create a new :obj:`~graph.GraphBuilder` object.
 
         Returns
         -------
-        :obj:`~_graph.GraphBuilder`
+        :obj:`~graph.GraphBuilder`
             Newly created graph builder object.
 
         """
+        from cuda.core.graph._graph_builder import GraphBuilder
+
         self._check_context_initialized()
         return GraphBuilder._init(stream=self.create_stream(), is_stream_owner=True)
 

cuda_core/cuda/core/_graph/__init__.py

@@ -26,7 +26,7 @@ def launch(stream: Stream | GraphBuilder | IsStreamT, config: LaunchConfig, kern
 
     Parameters
     ----------
-    stream : :obj:`~_stream.Stream` | :obj:`~_graph.GraphBuilder`
+    stream : :obj:`~_stream.Stream` | :obj:`~graph.GraphBuilder`
         The stream establishing the stream ordering semantic of a
         launch.
     config : :obj:`LaunchConfig`

cuda_core/cuda/core/_graph/_graph_def/__init__.py

@@ -182,7 +182,7 @@ cdef class Buffer:
 
         Parameters
         ----------
-        stream : :obj:`~_stream.Stream` | :obj:`~_graph.GraphBuilder`, optional
+        stream : :obj:`~_stream.Stream` | :obj:`~graph.GraphBuilder`, optional
             The stream object to use for asynchronous deallocation. If None,
             the deallocation stream stored in the handle is used.
         """
@@ -206,7 +206,7 @@ cdef class Buffer:
         ----------
         dst : :obj:`~_memory.Buffer`
             Source buffer to copy data from
-        stream : :obj:`~_stream.Stream` | :obj:`~_graph.GraphBuilder`
+        stream : :obj:`~_stream.Stream` | :obj:`~graph.GraphBuilder`
             Keyword argument specifying the stream for the
             asynchronous copy
 
@@ -237,7 +237,7 @@ cdef class Buffer:
         ----------
         src : :obj:`~_memory.Buffer`
             Source buffer to copy data from
-        stream : :obj:`~_stream.Stream` | :obj:`~_graph.GraphBuilder`
+        stream : :obj:`~_stream.Stream` | :obj:`~graph.GraphBuilder`
             Keyword argument specifying the stream for the
             asynchronous copy
 
@@ -262,7 +262,7 @@ cdef class Buffer:
         value : int | :obj:`collections.abc.Buffer`
             - int: Must be in range [0, 256). Converted to 1 byte.
             - :obj:`collections.abc.Buffer`: Must be 1, 2, or 4 bytes.
-        stream : :obj:`~_stream.Stream` | :obj:`~_graph.GraphBuilder`
+        stream : :obj:`~_stream.Stream` | :obj:`~graph.GraphBuilder`
             Stream for the asynchronous fill operation.
 
         Raises
@@ -496,7 +496,7 @@ cdef class MemoryResource:
         ----------
         size : int
             The size of the buffer to allocate, in bytes.
-        stream : :obj:`~_stream.Stream` | :obj:`~_graph.GraphBuilder`, optional
+        stream : :obj:`~_stream.Stream` | :obj:`~graph.GraphBuilder`, optional
             The stream on which to perform the allocation asynchronously.
             If None, it is up to each memory resource implementation to decide
             and document the behavior.
@@ -518,7 +518,7 @@ cdef class MemoryResource:
             The pointer or handle to the buffer to deallocate.
         size : int
             The size of the buffer to deallocate, in bytes.
-        stream : :obj:`~_stream.Stream` | :obj:`~_graph.GraphBuilder`, optional
+        stream : :obj:`~_stream.Stream` | :obj:`~graph.GraphBuilder`, optional
             The stream on which to perform the deallocation asynchronously.
             If None, it is up to each memory resource implementation to decide
             and document the behavior.

cuda_core/cuda/core/_launcher.pyx

@@ -129,7 +129,7 @@ cdef class _MemPool(MemoryResource):
         ----------
         size : int
             The size of the buffer to allocate, in bytes.
-        stream : :obj:`~_stream.Stream` | :obj:`~_graph.GraphBuilder`, optional
+        stream : :obj:`~_stream.Stream` | :obj:`~graph.GraphBuilder`, optional
             The stream on which to perform the allocation asynchronously.
             If None, an internal stream is used.
 
@@ -153,7 +153,7 @@ cdef class _MemPool(MemoryResource):
             The pointer or handle to the buffer to deallocate.
         size : int
             The size of the buffer to deallocate, in bytes.
-        stream : :obj:`~_stream.Stream` | :obj:`~_graph.GraphBuilder`, optional
+        stream : :obj:`~_stream.Stream` | :obj:`~graph.GraphBuilder`, optional
             The stream on which to perform the deallocation asynchronously.
             If the buffer is deallocated without an explicit stream, the allocation stream
             is used.

cuda_core/cuda/core/_memory/_buffer.pyx

@@ -38,7 +38,6 @@ from cuda.core._resource_handles cimport (
     as_py,
 )
 
-from cuda.core._graph import GraphBuilder
 
 
 @dataclass
@@ -354,17 +353,19 @@ cdef class Stream:
 
         return Stream._init(obj=_stream_holder())
 
-    def create_graph_builder(self) -> GraphBuilder:
-        """Create a new :obj:`~_graph.GraphBuilder` object.
+    def create_graph_builder(self) -> "GraphBuilder":
+        """Create a new :obj:`~graph.GraphBuilder` object.
 
         The new graph builder will be associated with this stream.
 
         Returns
         -------
-        :obj:`~_graph.GraphBuilder`
+        :obj:`~graph.GraphBuilder`
             Newly created graph builder object.
 
         """
+        from cuda.core.graph._graph_builder import GraphBuilder
+
         return GraphBuilder._init(stream=self, is_stream_owner=False)
 
 
@@ -466,6 +467,8 @@ cdef cydriver.CUstream _handle_from_stream_protocol(obj) except*:
 # Helper for API functions that accept either Stream or GraphBuilder. Performs
 # needed checks and returns the relevant stream.
 cdef Stream Stream_accept(arg, bint allow_stream_protocol=False):
+    from cuda.core.graph._graph_builder import GraphBuilder
+
     if isinstance(arg, Stream):
         return <Stream>(arg)
     elif isinstance(arg, GraphBuilder):

cuda_core/cuda/core/_memory/_memory_pool.pyx

@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2024-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 #
 # SPDX-License-Identifier: Apache-2.0
 
@@ -46,12 +46,6 @@ def _warn_deprecated():
 
 from cuda.core._device import Device
 from cuda.core._event import Event, EventOptions
-from cuda.core._graph import (
-    Graph,
-    GraphBuilder,
-    GraphCompleteOptions,
-    GraphDebugPrintOptions,
-)
 from cuda.core._launch_config import LaunchConfig
 from cuda.core._launcher import launch
 from cuda.core._layout import _StridedLayout
@@ -73,3 +67,9 @@ def _warn_deprecated():
 from cuda.core._module import Kernel, ObjectCode
 from cuda.core._program import Program, ProgramOptions
 from cuda.core._stream import Stream, StreamOptions
+from cuda.core.graph import (
+    Graph,
+    GraphBuilder,
+    GraphCompleteOptions,
+    GraphDebugPrintOptions,
+)

cuda_core/cuda/core/_stream.pyx

@@ -2,9 +2,9 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
-from cuda.core._graph._graph_def._graph_def cimport Condition, GraphDef
-from cuda.core._graph._graph_def._graph_node cimport GraphNode
-from cuda.core._graph._graph_def._subclasses cimport (
+from cuda.core.graph._graph_def cimport Condition, GraphDef
+from cuda.core.graph._graph_node cimport GraphNode
+from cuda.core.graph._subclasses cimport (
     AllocNode,
     ChildGraphNode,
     ConditionalNode,