Normalize the variant pack once, into a C type that is also the DLPack producer

dlpack_version.txt

@@ -1 +1 @@
-1.1
+1.3

docs/adding_torch_custom_ops.md

@@ -77,25 +77,25 @@ graph.check_support()
 graph.build_plans()  # also prepares variant pack template
 ```
 
-## Execute: use sorted_ptrs path
+## Execute
 
-Cache `uid_order` alongside the graph. Use `_execute_with_ptrs` instead of dict-based execute:
+`graph.execute(uid_to_tensor, workspace, handle=handle)` is the only form you need.
+It already caches the backend's operand order and reuses one pointer array, so the
+sorted-pointer path is what runs underneath — there is nothing faster to reach for,
+and hand-rolling it costs you the dynamic-shape overrides and the python-engine
+dispatch that `execute()` handles.
 
 ```python
 if cache_key not in _cache:
-    graph, ws_size = _build_graph(...)
-    uid_order = graph._get_variant_pack_uids_sorted()
-    _cache[cache_key] = (graph, ws_size, uid_order)
+    _cache[cache_key] = _build_graph(...)          # (graph, ws_size)
 
-graph, ws_size, uid_order = _cache[cache_key]
+graph, ws_size = _cache[cache_key]
 
 # Allocate workspace per-call (PyTorch's caching allocator recycles efficiently)
 workspace = torch.empty(max(ws_size, 1), device=x.device, dtype=torch.uint8)
 
-# Build uid→tensor map, extract sorted ptrs
 uid_to_tensor = {X.get_uid(): x, W.get_uid(): w, Y.get_uid(): y_out}
-ptrs = [uid_to_tensor[uid].data_ptr() for uid in uid_order]
-graph._execute_with_ptrs(ptrs, workspace.data_ptr(), int(handle))
+graph.execute(uid_to_tensor, workspace, handle=handle)
 ```
 
 **Do NOT cache workspace tensors** — they can race on different CUDA streams.
@@ -175,7 +175,7 @@ def my_op(x, w, eps=1e-5, bias=None):
 
 - [ ] Use `torch.Library.define/impl`, not `@torch.library.custom_op`
 - [ ] Cache graph + uid_order + workspace in module-level dict
-- [ ] Use `graph._execute_with_ptrs(sorted_ptrs)` not `graph.execute(dict)`
+- [ ] Use `graph.execute(uid_to_tensor, workspace, handle=handle)` — it takes the sorted-pointer path internally
 - [ ] Use explicit UIDs (IntEnum) for stable cache keys
 - [ ] Cache cuDNN handle per device
 - [ ] Allocate workspace per-call (do NOT cache — stream safety)
@@ -192,7 +192,7 @@ def my_op(x, w, eps=1e-5, bias=None):
 | `torch.empty` per output tensor | ~2 each | consider `out=` or pre-alloc |
 | cache key build + lookup | ~1.5 | tuple construction + dict hash |
 | uid→tensor dict + list comp | ~1 | Python overhead |
-| `graph._execute_with_ptrs` | ~19 | 1.7 us FE + 0.8 us varpack + 5.6 us backend |
+| `graph.execute` | ~19 | 1.7 us FE + 0.8 us varpack + 5.6 us backend |
 | **Total (well-optimized)** | **~52** | vs ~18 us for native ATen ops |
 
 The ~34 us gap vs native ATen is torch.ops dispatcher + autograd overhead.

docs/python_graph_and_execution_backends.md

@@ -80,15 +80,65 @@ create_execution_plans([heur_mode.A, ...])                    _pygraph.py
 - `BaseEngine`: `check_support(graph)` (accept, or decline by raising),
   `build_plan(graph, plan, ctx) → CompiledPlan` (the expensive
   JIT step, once per graph/plan, cached on the graph),
-  `CompiledPlan.execute(graph, uid_to_data, ExecutionContext)` with explicit
-  handle/stream/workspace/overrides. `uid_to_data` is the caller's variant
-  pack, exactly as the classic backend receives it; engines that address
-  buffers by port name call `resolve_node_buffers(graph, uid_to_data)`
-  (`engines/base.py`), which joins the pack with each node's wired ports —
-  strict missing-buffer validation, torch tensors detached once (DLPack/CAI
-  refuse `requires_grad` export) — into per-node `NodeBuffers`
-  (`{port_name: caller buffer}`). Simple eager engines implement
-  `execute()` only.
+  `CompiledPlan.execute(graph, operands, ExecutionContext)` with explicit
+  handle/stream/workspace. Dynamic-shape overrides are a backend-path feature:
+  a python plan is compiled for the shapes the graph declared, so `execute()`
+  refuses them rather than silently running a different problem. Simple eager
+  engines implement `execute()` only.
+
+#### The variant pack is normalized once
+
+`graph.execute()` converts whatever the caller passed — a torch tensor, a
+`DeviceView`, any `__dlpack__` / `__cuda_array_interface__` producer, or a bare
+device address — into a `VariantPack` at the top, and everything below reads that.
+**Do not add a branch on what the caller's object is.** This exists because
+there used to be two such branches: the backend path accepted a bare address
+(`_native_var_pack`'s `if type(d) is int: return d`) while an engine got the
+object untouched and `frost.buffers.probe` refused it. One public call, two
+answers, decided by which plan the heuristics happened to pick — which the
+caller does not control.
+
+`VariantPack` carries the caller-filled uids ascending and the operands
+themselves, held in a C type (`pygraph/variant_pack.cpp`) as one `DLTensor`
+each. That type both consumes `__dlpack_c_exchange_api__` — the C function
+table a producer publishes on its type, which is how one crossing reads the
+whole pack — and implements it, so a kernel reads a slot through the same fast
+path it has for a framework tensor. `address` is the pointer array
+`_execute_with_raw_ptrs` takes.
+
+Each operand's OWN dim/stride/data_type is what the pack holds, deliberately
+not the graph's declaration: the two may differ and one engine relies on it —
+`frost_gemm` takes its M/N/K from the buffers, so a plan built for one problem
+size runs another bit-exactly. **Read the IR port for the shape the plan was
+built for; read the pack for the shape about to run.**
+
+Two rules that are easy to break by accident:
+
+- **The pointer array is per call.** Two threads may execute one graph
+  concurrently with different buffers. A shared array hands each thread the
+  other's pointers, and each pointer in it is individually valid, so the failure
+  is a wrong number rather than a raise.
+- **The operand order has exactly one source, never a union.** The lowered
+  graph's variant-pack template when the graph has one — only C++ sees every
+  user slot, since a tensor's `ragged_offset` is an operand but hangs off the
+  `Tensor` rather than off a node port, and the slots the graph fills itself
+  (pass-by-value scalars, slice replacement destinations, workspace
+  modifications) must be excluded. The python IR only for the python-only ops
+  that never lower. The two sides do not have to agree: each indexes the layout
+  it was handed.
+
+`is_virtual` on a python-only graph does not mean "the caller supplies
+nothing" — it is a statement about the backend's lowering, and a gdn graph marks
+its own `O` virtual while the caller passes a buffer for it. So the layout there
+is every wired port, and an unfilled slot is an optional port the caller did not
+request.
+
+`CompiledPlan.takes_variant_pack` is the migration flag. An engine that has not
+set it still receives the caller's `{uid: buffer}` map and reaches ports through
+`resolve_node_buffers`; the flag and that function both go once the last engine
+has moved. `execute()` builds the `Tensor`s only for a plan that sets the flag —
+measured, normalizing for a plan that will not read the result costs more than
+it saves.
 - **An engine does not propose its own plans.** Which configs to try, in what
   order, and where the backend's entries belong is one comparison across every
   candidate, and no engine can make it from the inside — it sees neither its

python/CMakeLists.txt

@@ -13,11 +13,16 @@ set(BUILD_MOCK OFF)
 option(CUDNN_FRONTEND_USE_SYSTEM_DLPACK "Whether dlpack should use the system version or fetched." OFF)
 
 IF(CUDNN_FRONTEND_USE_SYSTEM_DLPACK)
-  find_package(dlpack REQUIRED)
+  # 1.3 is where DLPackExchangeAPI is declared, which pygraph/variant_pack.cpp
+  # needs. The wire structs are byte-identical to 1.1 -- sizeof and every
+  # offsetof of DLTensor and DLManagedTensor match -- so this is a
+  # compile-time requirement only; capsules exchanged with a consumer built
+  # against 1.1 are unaffected.
+  find_package(dlpack 1.3 REQUIRED)
   if(dlpack_FOUND)
-    message(STATUS "Found system dlpack")
+    message(STATUS "Found system dlpack ${dlpack_VERSION}")
   else()
-    message(FATAL_ERROR "dlpack not found")
+    message(FATAL_ERROR "dlpack >= 1.3 not found (needed for DLPackExchangeAPI); unset CUDNN_FRONTEND_USE_SYSTEM_DLPACK to fetch it")
   endif()
 else()
 
@@ -68,6 +73,7 @@ python_add_library(
     pygraph/norm.cpp
     pygraph/sdpa.cpp
     pygraph/pointwise.cpp
+    pygraph/variant_pack.cpp
 
     WITH_SOABI
 )

python/cudnn/__init__.py

@@ -120,93 +120,6 @@ def _set_data_type(
 _pybind_module.backend_graph.tensor = _tensor
 
 
-def _library_device_pointer(input_tensor):
-    # either pass in pointers directly
-    if type(input_tensor) is int:
-        return input_tensor
-    # directly extract data pointer for torch tensors
-    elif _is_torch_tensor(input_tensor):
-        return input_tensor.data_ptr()
-    # fall back to dlpack support by library
-    else:
-        return _pybind_module._get_data_ptr(input_tensor)
-
-
-def _execute(
-    self,
-    tensor_to_device_buffer,
-    workspace,
-    handle=None,
-    override_uids=None,
-    override_shapes=None,
-    override_strides=None,
-):
-    """
-    Execute a cudnn graph.
-
-    Args:
-        tensor_to_device_buffer (dict(cudnn_tensor, Union[torch.Tensor, int, __dlpack__])): The dimensions of the tensor.
-        workspace (Union[torch.Tensor, int, __dlpack__]): The name of the tensor.
-        handle: cudnn_handle created with cudnn.create_handle()
-    Returns:
-        None
-    """
-    uid_to_tensor_pointer = {
-        x if type(x) is int else x.get_uid(): _library_device_pointer(pointer) for x, pointer in tensor_to_device_buffer.items() if x is not None
-    }
-
-    workspace_pointer = _library_device_pointer(workspace)
-    self._execute(
-        uid_to_tensor_pointer,
-        workspace_pointer,
-        handle,
-        override_uids,
-        override_shapes,
-        override_strides,
-    )
-
-
-def _execute_plan_at_index(
-    self,
-    tensor_to_device_buffer,
-    workspace,
-    index,
-    handle=None,
-    override_uids=None,
-    override_shapes=None,
-    override_strides=None,
-):
-    """
-    Execute a cudnn graph.
-
-    Args:
-        tensor_to_device_buffer (dict(cudnn_tensor, Union[torch.Tensor, int, __dlpack__])): The dimensions of the tensor.
-        workspace (Union[torch.Tensor, int, __dlpack__]): The name of the tensor.
-        index(int): Location of execution plan to use.
-        handle: cudnn_handle created with cudnn.create_handle()
-    Returns:
-        None
-    """
-    uid_to_tensor_pointer = {
-        x if type(x) is int else x.get_uid(): _library_device_pointer(pointer) for x, pointer in tensor_to_device_buffer.items() if x is not None
-    }
-
-    workspace_pointer = _library_device_pointer(workspace)
-    self._execute_plan_at_index(
-        uid_to_tensor_pointer,
-        workspace_pointer,
-        index,
-        handle,
-        override_uids,
-        override_shapes,
-        override_strides,
-    )
-
-
-_pybind_module.backend_graph.execute = _execute
-_pybind_module.backend_graph.execute_plan_at_index = _execute_plan_at_index
-
-
 def load_cudnn():
     # First look at python site packages
     lib_path = glob.glob(os.path.join(sysconfig.get_path("purelib"), "nvidia/cudnn/bin/cudnn64_9.dll"))