Skip to content

Commit 0aa57e0

Browse files
committed
Stage F: ExecutableGraph runtime executor over auto-compiled fusion regions
Stages A-E in cppmega_v4.fusion plan regions, build the V4-extended descriptor registry, and emit AutoCompiledRegion records, but nothing walks those regions and runs them — auto_compile_plan is a list of intents until something dispatches them. This commit adds that missing consumer. - cppmega_v4/fusion/execute.py (new) - ExecutableGraph: holds BrickGraph + planner plans + compiled regions; per-region artifact slot (Callable[[mx.array], mx.array]); forward(hidden) walks regions in planner order, dispatching to the attached artifact when present or to the bricks' native nn.Module forward otherwise. Per-region timing + backend choice + canonical eager-reason string land in execution_log on every call so the GUI/matrix can prove which regions actually fused. - RegionExecution: per-region outcome record (backend, brick names, eager_reason, duration_ns). - build_executable_graph(model, ...): convenience wrapper that plans regions via plan_fusion_regions + auto_compile_plan and returns an ExecutableGraph with no artifacts attached yet. - cppmega_v4/fusion/__init__.py: surface ExecutableGraph, RegionExecution, build_executable_graph. - tests/v4/test_fusion_stage_f.py: 11 cases covering eager forward parity, region-summary shape, execution-log labelling for single- brick and dlpack-handoff regions, attach_artifact happy path + rejection cases (non-callable, out-of-range, passthrough region), and the explicit instantiation-required failure mode for cheap block-spec graphs. Forward-only by design — training (value_and_grad, backward) keeps going through the m04 fused train-block runtime; nothing in this seam allocates banks or copies parameter tensors. Eager-region fallback keeps the no-hidden-allocation rule intact. Verified: pytest -q on the full v4 fusion suite (Stage A-E + the new Stage F + roadmap gaps + the Path C physical ABI bank owner test) -> 405 passed; ruff and pyright clean on the new module.
1 parent 91000dc commit 0aa57e0

3 files changed

Lines changed: 578 additions & 0 deletions

File tree

cppmega_v4/fusion/__init__.py

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -52,6 +52,11 @@
5252
mlx_to_tilelang,
5353
tilelang_to_mlx,
5454
)
55+
from cppmega_v4.fusion.execute import (
56+
ExecutableGraph,
57+
RegionExecution,
58+
build_executable_graph,
59+
)
5560

5661
__all__ = [
5762
"DEFAULT_MAX_REGION_SIZE",
@@ -68,6 +73,9 @@
6873
"auto_fuse_model",
6974
"build_v4_extended_registry",
7075
"can_fuse_pair",
76+
"ExecutableGraph",
77+
"RegionExecution",
78+
"build_executable_graph",
7179
"detect_region_pattern",
7280
"dlpack_available",
7381
"from_block_specs",

cppmega_v4/fusion/execute.py

Lines changed: 325 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,325 @@
1+
"""Stage F — executable fusion graph over auto-compiled brick regions.
2+
3+
This module provides the runtime executor that downstream consumers were
4+
missing in Stages A-E:
5+
6+
* Stages A-E plan and *describe* fusion regions, register V4 brick
7+
descriptors with the cppmega_mlx schedule machinery, and emit
8+
AutoCompiledRegion records that pair a planner region with the
9+
TileLang schedule template the codegen consumes.
10+
* Stage F walks those regions in their planner order and executes the
11+
contained bricks against MLX inputs, in the right order, with the
12+
backend the planner chose:
13+
- ``backend == "path_c"`` and the region has a compiled artifact
14+
attached: drive the artifact's forward (artifact must accept the
15+
region's leading hidden state via DLPack and write its tail to
16+
the same buffer convention).
17+
- otherwise (single-brick passthrough, dlpack-handoff chain,
18+
metal_inline, or any region whose Path C lowering has not yet
19+
landed): execute every brick eagerly via its existing
20+
``nn.Module`` instance, threading the running hidden state from
21+
one brick to the next.
22+
23+
The executor is intentionally honest about which regions are running
24+
fused vs eager — see :class:`RegionExecution.backend` and the
25+
:attr:`ExecutableGraph.execution_log` field — so callers (1B path_c
26+
matrix, GUI, benchmarks) can decide whether the planner's promises are
27+
being honoured for a given input. No tensor data is silently copied or
28+
re-typed beyond the explicit DLPack handoff helpers in
29+
:mod:`cppmega_v4.fusion.dlpack_bridge`.
30+
31+
Public surface:
32+
- :class:`ExecutableGraph` — bound to one mlx model + one plan.
33+
- :class:`RegionExecution` — per-region outcome record.
34+
- :func:`build_executable_graph(model, plan_regions=None)` —
35+
convenience entry point that auto-plans the regions from
36+
``auto_fuse_model`` when ``plan_regions`` is omitted.
37+
"""
38+
39+
from __future__ import annotations
40+
41+
from collections.abc import Callable, Mapping, Sequence
42+
from dataclasses import dataclass, field
43+
import time
44+
from typing import Any
45+
46+
import mlx.core as mx
47+
import mlx.nn as nn
48+
49+
from cppmega_v4.fusion.auto_compile import (
50+
AutoCompiledRegion,
51+
RegionPattern,
52+
auto_compile_plan,
53+
)
54+
from cppmega_v4.fusion.auto_planner import (
55+
DEFAULT_MAX_REGION_SIZE,
56+
DEFAULT_MAX_SHARED_MEM_BYTES,
57+
FusionRegionPlan,
58+
plan_fusion_regions,
59+
)
60+
from cppmega_v4.fusion.brick_graph import (
61+
BrickGraph,
62+
BrickNode,
63+
from_mlx_model,
64+
)
65+
66+
67+
# Reasons a region must run via the eager-brick path instead of a fused
68+
# artifact. Keep these stable — the matrix/GUI surface them verbatim.
69+
_REGION_RUN_EAGER_REASONS: dict[str, str] = {
70+
"single_brick": "single-brick region; no fused PrimFunc emitted",
71+
"dlpack_handoff": "planner chose dlpack_handoff; each brick keeps its native kernel",
72+
"no_template": "AutoCompiledRegion has no schedule template",
73+
"no_artifact": "no compiled TileLang artifact attached to this region",
74+
"no_module": "brick instance missing module reference (instantiate=True required)",
75+
}
76+
77+
78+
@dataclass(frozen=True)
79+
class RegionExecution:
80+
"""One region's outcome from a single ExecutableGraph.forward call.
81+
82+
Fields:
83+
region_index: position in the planner output, 0-based.
84+
pattern: planner pattern label (single-brick, fused, etc.).
85+
backend: "path_c_artifact" | "eager_bricks" — what actually ran.
86+
brick_names: ordered tuple of brick names inside this region.
87+
eager_reason: when ``backend == "eager_bricks"``, the canonical
88+
reason string from :data:`_REGION_RUN_EAGER_REASONS` (or
89+
free-form fallback). ``""`` when the artifact ran.
90+
duration_ns: wall-clock spent in this region's execution call.
91+
"""
92+
93+
region_index: int
94+
pattern: RegionPattern
95+
backend: str
96+
brick_names: tuple[str, ...]
97+
eager_reason: str
98+
duration_ns: int
99+
100+
@property
101+
def ran_fused(self) -> bool:
102+
return self.backend == "path_c_artifact"
103+
104+
105+
@dataclass
106+
class ExecutableGraph:
107+
"""Forward-only executor for a BrickGraph + planner regions + bricks.
108+
109+
The executor never owns the bricks — it dispatches to their existing
110+
``nn.Module`` instances. Per-region compiled artifacts can be
111+
attached after construction via :meth:`attach_artifact` and will be
112+
used for that region's forward pass when present. Otherwise the
113+
region falls back to walking the bricks eagerly with a sequential
114+
hidden-state thread.
115+
116+
The executor is intentionally forward-only. Training (value_and_grad,
117+
backward) goes through the m04 fused train-block runtime, which is a
118+
separate seam — see ``scripts/m04_train_step.py``.
119+
"""
120+
121+
graph: BrickGraph
122+
plans: tuple[FusionRegionPlan, ...]
123+
regions: tuple[AutoCompiledRegion, ...]
124+
# Mutable: per-region attached artifact (callable; takes leading
125+
# hidden mx.array, returns trailing hidden mx.array). Populated by
126+
# callers via attach_artifact after compile.
127+
artifacts: dict[int, Callable[[mx.array], mx.array]] = field(
128+
default_factory=dict
129+
)
130+
# Mutable: rolling per-region execution log (last forward call).
131+
execution_log: list[RegionExecution] = field(default_factory=list)
132+
133+
# --- introspection -------------------------------------------------
134+
135+
@property
136+
def fused_region_count(self) -> int:
137+
return sum(1 for region in self.regions if region.has_compiled_template)
138+
139+
@property
140+
def eager_region_count(self) -> int:
141+
return len(self.regions) - self.fused_region_count
142+
143+
def region_summary(self) -> tuple[dict[str, Any], ...]:
144+
"""Per-region summary suitable for receipts / GUI tables."""
145+
146+
return tuple(
147+
{
148+
"index": index,
149+
"pattern": region.pattern.value,
150+
"backend_planned": region.plan.backend,
151+
"brick_names": list(region.plan.brick_names),
152+
"has_compiled_template": region.has_compiled_template,
153+
"has_artifact": index in self.artifacts,
154+
"estimated_savings_us": float(region.plan.estimated_savings_us),
155+
}
156+
for index, region in enumerate(self.regions)
157+
)
158+
159+
# --- artifact wiring -----------------------------------------------
160+
161+
def attach_artifact(
162+
self,
163+
region_index: int,
164+
artifact: Callable[[mx.array], mx.array],
165+
) -> None:
166+
"""Bind a compiled TileLang artifact to a specific region index.
167+
168+
The artifact must accept the leading hidden state as its single
169+
positional ``mx.array`` argument and return the trailing hidden
170+
state as an ``mx.array``. Any device or layout requirements are
171+
the artifact's responsibility — callers are expected to wire
172+
DLPack-friendly artifacts produced by Path C codegen.
173+
"""
174+
if not callable(artifact):
175+
raise TypeError(
176+
f"artifact for region {region_index} must be callable"
177+
)
178+
if region_index < 0 or region_index >= len(self.regions):
179+
raise IndexError(
180+
f"region_index {region_index} out of range "
181+
f"[0, {len(self.regions)})"
182+
)
183+
region = self.regions[region_index]
184+
if not region.has_compiled_template:
185+
raise ValueError(
186+
f"region {region_index} ({region.pattern.value}) has no "
187+
"compiled schedule template; cannot attach a fused artifact"
188+
)
189+
self.artifacts[region_index] = artifact
190+
191+
# --- forward -------------------------------------------------------
192+
193+
def forward(self, hidden: mx.array) -> mx.array:
194+
"""Run every region in planner order against the input ``hidden``.
195+
196+
Each region's leading brick consumes the current hidden state;
197+
the trailing brick's output becomes the next region's input.
198+
Per-region timing and backend choice are recorded in
199+
:attr:`execution_log` (replaced on every call).
200+
"""
201+
self.execution_log.clear()
202+
running = hidden
203+
for index, region in enumerate(self.regions):
204+
t0 = time.perf_counter_ns()
205+
artifact = self.artifacts.get(index)
206+
if artifact is not None:
207+
running = artifact(running)
208+
backend = "path_c_artifact"
209+
eager_reason = ""
210+
else:
211+
running = self._run_region_eager(region, running)
212+
backend = "eager_bricks"
213+
eager_reason = self._eager_reason_for(region)
214+
duration = time.perf_counter_ns() - t0
215+
self.execution_log.append(
216+
RegionExecution(
217+
region_index=index,
218+
pattern=region.pattern,
219+
backend=backend,
220+
brick_names=region.plan.brick_names,
221+
eager_reason=eager_reason,
222+
duration_ns=int(duration),
223+
)
224+
)
225+
return running
226+
227+
__call__ = forward
228+
229+
# --- helpers -------------------------------------------------------
230+
231+
def _run_region_eager(
232+
self,
233+
region: AutoCompiledRegion,
234+
hidden: mx.array,
235+
) -> mx.array:
236+
"""Walk a region's bricks eagerly, threading the hidden state."""
237+
running = hidden
238+
for brick_name in region.plan.brick_names:
239+
node = self._node_by_name(brick_name)
240+
module = node.module
241+
if module is None:
242+
raise ValueError(
243+
f"brick {brick_name!r} has no module attached; the "
244+
"executor needs instantiated modules to run eagerly. "
245+
"Re-build the graph with instantiate=True or attach "
246+
"the module via BrickNode.module."
247+
)
248+
out = module(running)
249+
running = self._sanitise_brick_output(out, running, brick_name)
250+
return running
251+
252+
def _node_by_name(self, name: str) -> BrickNode:
253+
for node in self.graph.nodes:
254+
if node.name == name:
255+
return node
256+
raise KeyError(name)
257+
258+
@staticmethod
259+
def _sanitise_brick_output(
260+
out: Any, prev_hidden: mx.array, brick_name: str
261+
) -> mx.array:
262+
"""Brick modules sometimes return tuples ``(hidden, state)``.
263+
264+
Reduce them to the leading hidden ``mx.array`` so the next brick
265+
sees a clean state. Returns ``prev_hidden`` if the brick output
266+
is not a usable ``mx.array``.
267+
"""
268+
if isinstance(out, mx.array):
269+
return out
270+
if isinstance(out, tuple) and out and isinstance(out[0], mx.array):
271+
return out[0]
272+
if isinstance(out, Mapping):
273+
for value in out.values():
274+
if isinstance(value, mx.array):
275+
return value
276+
return prev_hidden
277+
278+
def _eager_reason_for(self, region: AutoCompiledRegion) -> str:
279+
if region.pattern is RegionPattern.SINGLE_BRICK_PASSTHROUGH:
280+
return _REGION_RUN_EAGER_REASONS["single_brick"]
281+
if region.pattern is RegionPattern.DLPACK_HANDOFF_CHAIN:
282+
return _REGION_RUN_EAGER_REASONS["dlpack_handoff"]
283+
if not region.has_compiled_template:
284+
return _REGION_RUN_EAGER_REASONS["no_template"]
285+
return _REGION_RUN_EAGER_REASONS["no_artifact"]
286+
287+
288+
# ---------------------------------------------------------------------------
289+
# Convenience entry point
290+
# ---------------------------------------------------------------------------
291+
292+
293+
def build_executable_graph(
294+
model: nn.Module,
295+
*,
296+
attr_order: Sequence[str] | None = None,
297+
max_region_size: int = DEFAULT_MAX_REGION_SIZE,
298+
max_shared_mem_bytes: int = DEFAULT_MAX_SHARED_MEM_BYTES,
299+
) -> ExecutableGraph:
300+
"""Walk ``model``, plan regions, auto-compile descriptors, and bind.
301+
302+
Returns an :class:`ExecutableGraph` with no artifacts attached yet —
303+
every region defaults to eager execution. Callers wire artifacts
304+
explicitly via :meth:`ExecutableGraph.attach_artifact` after their
305+
own compile step (so we never mix the planner with TileLang codegen
306+
here).
307+
"""
308+
graph = from_mlx_model(model, attr_order=attr_order)
309+
plans = tuple(
310+
plan_fusion_regions(
311+
graph,
312+
max_region_size=max_region_size,
313+
max_shared_mem_bytes=max_shared_mem_bytes,
314+
)
315+
)
316+
kinds_by_name = {n.name: n.kind for n in graph.nodes}
317+
regions = tuple(auto_compile_plan(list(plans), kinds_by_name))
318+
return ExecutableGraph(graph=graph, plans=plans, regions=regions)
319+
320+
321+
__all__ = [
322+
"ExecutableGraph",
323+
"RegionExecution",
324+
"build_executable_graph",
325+
]

0 commit comments

Comments
 (0)