Skip to content

Commit 55066d8

Browse files
committed
feat(v4/fusion): Stage A — brick_graph + dlpack_bridge + compatibility
First step of the Auto-Fusion roadmap (Auto-FusionLayerBricks.md): foundation layer that downstream stages (descriptor synth, auto-planner, architecture presets, in-region auto-compile) all consume. - cppmega_v4/fusion/brick_graph.py: BrickNode / BrickGraph + two ingest paths: from_block_specs (JSON-shaped GUI surface) and from_mlx_model (walks nn.Module tree, honours ``_v4_brick_kind`` instance marker for closure-defined builders). Validates kind against BLOCK_BUILDERS, rejects duplicate names and dangling edges eagerly. - cppmega_v4/fusion/dlpack_bridge.py: mlx_to_tilelang / tilelang_to_mlx zero-copy via DLPack (Metal-resident capsules per mlx#2848 PoC), dlpack_available() probe (lru_cached), host_copy_fallback for the last-resort path. - cppmega_v4/fusion/compatibility.py: table-driven FusionEligibility oracle. Per-kind categories (linear_attn / sdpa_attention / ssm / nonlinear_rnn / moe / cross_attn / sparse_attn / norm_or_proj) + symmetric pair rules with explicit backend ("path_c" / "metal_inline" / "dlpack_handoff") and human reason. Hard boundaries (MoE routing, sparse_attn) enforced. System tests verify the oracle correctly classifies real model patterns: Qwen3-Next (3 GDN fuse, attn boundary, MoE boundary) and Ling-2.6 (7 bailing_linear fuse, MLA boundary). Tests: 24/24 new pass; full v4 suite 422 passed / 5 skipped / 0 failed.
1 parent 00cdaf0 commit 55066d8

6 files changed

Lines changed: 1154 additions & 0 deletions

File tree

Auto-FusionLayerBricks.md

Lines changed: 238 additions & 0 deletions
Large diffs are not rendered by default.

cppmega_v4/fusion/__init__.py

Lines changed: 48 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,48 @@
1+
"""Auto-fusion layer over V4 bricks.
2+
3+
See ``Auto-FusionLayerBricks.md`` (repo root) for the full design.
4+
5+
This package walks composed UnifiedSuperblock graphs of ``BLOCK_BUILDERS``
6+
bricks, detects fusion-eligible adjacencies via a table-driven oracle, and
7+
hands compiled regions to ``cppmega_mlx.runtime.path_c_fusion`` for
8+
TileLang lowering. Tensors cross fusion boundaries zero-copy via DLPack.
9+
10+
Stage A surface (this commit):
11+
- brick_graph.BrickNode / BrickGraph + walkers
12+
- dlpack_bridge.mlx_to_tilelang / tilelang_to_mlx
13+
- compatibility.FusionEligibility / can_fuse_pair
14+
15+
Stages B-E are TBD (see roadmap).
16+
"""
17+
18+
from __future__ import annotations
19+
20+
from cppmega_v4.fusion.brick_graph import (
21+
BrickGraph,
22+
BrickNode,
23+
from_block_specs,
24+
from_mlx_model,
25+
)
26+
from cppmega_v4.fusion.compatibility import (
27+
FusionEligibility,
28+
can_fuse_pair,
29+
)
30+
from cppmega_v4.fusion.dlpack_bridge import (
31+
dlpack_available,
32+
host_copy_fallback,
33+
mlx_to_tilelang,
34+
tilelang_to_mlx,
35+
)
36+
37+
__all__ = [
38+
"BrickGraph",
39+
"BrickNode",
40+
"FusionEligibility",
41+
"can_fuse_pair",
42+
"dlpack_available",
43+
"from_block_specs",
44+
"from_mlx_model",
45+
"host_copy_fallback",
46+
"mlx_to_tilelang",
47+
"tilelang_to_mlx",
48+
]

cppmega_v4/fusion/brick_graph.py

Lines changed: 237 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,237 @@
1+
"""BrickGraph — graph IR over composed V4 bricks.
2+
3+
A V4 model is a sequence (or DAG, for MoE routing) of bricks chosen from
4+
``cppmega_v4.models.unified_superblock_v4.BLOCK_BUILDERS``. BrickGraph is
5+
the smallest representation that downstream fusion planners need:
6+
7+
- what brick (kind + params + instantiated module)
8+
- producer→consumer dependencies (linear chain or routed fan-out/in)
9+
10+
Two ingest paths:
11+
12+
``from_block_specs(specs)`` — JSON-shaped block specs (the same surface
13+
the future GUI emits). Each spec is ``{"kind": "...", "name": "...",
14+
"params": {...}}``. Linear chain inferred from list order.
15+
16+
``from_mlx_model(model)`` — walk an existing ``nn.Module`` tree and
17+
identify which submodules are V4 bricks (by class match against
18+
BLOCK_BUILDERS' construction output). Linear chain inferred from
19+
attribute iteration order on the parent module.
20+
"""
21+
22+
from __future__ import annotations
23+
24+
from collections.abc import Sequence
25+
from dataclasses import dataclass, field
26+
from typing import Any
27+
28+
import mlx.nn as nn
29+
30+
from cppmega_v4.models.unified_superblock_v4 import BLOCK_BUILDERS
31+
32+
33+
_ADDITIONAL_FUSION_KINDS = frozenset(
34+
{
35+
"mamba3",
36+
"m2rnn",
37+
"sparse_mla_fp8",
38+
"ssm",
39+
}
40+
)
41+
42+
43+
@dataclass(frozen=True)
44+
class BrickNode:
45+
"""One brick in the graph: kind, name, params, instantiated module."""
46+
47+
kind: str
48+
name: str
49+
params: dict[str, Any] = field(default_factory=dict)
50+
module: nn.Module | None = None
51+
52+
def __post_init__(self) -> None:
53+
if not self.kind:
54+
raise ValueError("BrickNode.kind must be non-empty")
55+
if not self.name:
56+
raise ValueError("BrickNode.name must be non-empty")
57+
if self.kind not in BLOCK_BUILDERS and self.kind not in _ADDITIONAL_FUSION_KINDS:
58+
raise ValueError(
59+
f"BrickNode.kind={self.kind!r} not in BLOCK_BUILDERS "
60+
f"({sorted(BLOCK_BUILDERS)!r})"
61+
)
62+
63+
64+
@dataclass(frozen=True)
65+
class BrickGraph:
66+
"""Directed graph of bricks. Edges are ``(producer_name, consumer_name)``.
67+
68+
A simple linear chain has ``edges = [(n0,n1), (n1,n2), ...]``. MoE
69+
routing introduces fan-out: one router node with N outbound edges, one
70+
combine node with N inbound edges.
71+
"""
72+
73+
nodes: tuple[BrickNode, ...]
74+
edges: tuple[tuple[str, str], ...] = ()
75+
76+
def __post_init__(self) -> None:
77+
seen: set[str] = set()
78+
for n in self.nodes:
79+
if n.name in seen:
80+
raise ValueError(f"duplicate brick name {n.name!r}")
81+
seen.add(n.name)
82+
names = {n.name for n in self.nodes}
83+
for p, c in self.edges:
84+
if p not in names:
85+
raise ValueError(f"edge producer {p!r} not in node set")
86+
if c not in names:
87+
raise ValueError(f"edge consumer {c!r} not in node set")
88+
89+
@property
90+
def names(self) -> tuple[str, ...]:
91+
return tuple(n.name for n in self.nodes)
92+
93+
def by_name(self, name: str) -> BrickNode:
94+
for n in self.nodes:
95+
if n.name == name:
96+
return n
97+
raise KeyError(name)
98+
99+
def successors(self, name: str) -> tuple[str, ...]:
100+
return tuple(c for p, c in self.edges if p == name)
101+
102+
def predecessors(self, name: str) -> tuple[str, ...]:
103+
return tuple(p for p, c in self.edges if c == name)
104+
105+
106+
# ---------------------------------------------------------------------------
107+
# Ingest from JSON-shaped block specs
108+
# ---------------------------------------------------------------------------
109+
110+
111+
def from_block_specs(
112+
specs: Sequence[dict],
113+
*,
114+
hidden_size: int,
115+
instantiate: bool = True,
116+
) -> BrickGraph:
117+
"""Build a BrickGraph from a list of block specs.
118+
119+
Each spec:
120+
{"kind": "<BLOCK_BUILDERS key>", "name": "<unique>", "params": {...}}
121+
122+
Edges are inferred as a linear chain (specs[i] → specs[i+1]).
123+
124+
Args:
125+
specs: ordered list of brick specs.
126+
hidden_size: passed to each ``BLOCK_BUILDERS[kind](hidden_size, params)``.
127+
instantiate: when True, construct each module via BLOCK_BUILDERS.
128+
When False, BrickNode.module is None (cheap path for planners
129+
that only need kind/params).
130+
"""
131+
nodes: list[BrickNode] = []
132+
used_names: set[str] = set()
133+
for i, spec in enumerate(specs):
134+
kind = spec["kind"]
135+
name = spec.get("name") or f"{kind}_{i}"
136+
if name in used_names:
137+
raise ValueError(f"duplicate name {name!r} in specs[{i}]")
138+
used_names.add(name)
139+
params = dict(spec.get("params") or {})
140+
module = (
141+
BLOCK_BUILDERS[kind](hidden_size, params) if instantiate else None
142+
)
143+
nodes.append(BrickNode(kind=kind, name=name, params=params, module=module))
144+
145+
edges = tuple(
146+
(nodes[i].name, nodes[i + 1].name) for i in range(len(nodes) - 1)
147+
)
148+
return BrickGraph(nodes=tuple(nodes), edges=edges)
149+
150+
151+
# ---------------------------------------------------------------------------
152+
# Ingest from an existing nn.Module tree
153+
# ---------------------------------------------------------------------------
154+
155+
156+
def _brick_kind_of(module: nn.Module) -> str | None:
157+
"""Return the BLOCK_BUILDERS kind that produced ``module``, or None.
158+
159+
The check is structural: each builder returns a known class (or wraps
160+
upstream in our V4 wrapper). We look up by qualified class name.
161+
"""
162+
cls_name = type(module).__name__
163+
# Hard-coded map of class-name -> BLOCK_BUILDERS key. Mirrors the
164+
# builders in cppmega_v4/models/unified_superblock_v4.py. New brick
165+
# added to BLOCK_BUILDERS must also be added here for graph walking.
166+
KIND_BY_CLASS: dict[str, str] = {
167+
"GatedAttentionBlock": "gated_attention",
168+
"Mistral4MLABlock": "mistral4_mla",
169+
"DSv4AttentionBlock": "dsv4_attention",
170+
"BailingLinearAttnBlock": "bailing_linear",
171+
"BailingMLABlock": "bailing_mla",
172+
"BailingMoEBlock": "bailing_moe",
173+
"Gemma4DrafterLayerBlock": "gemma4_drafter",
174+
"NemotronHMTPBlockWrapper": "nemotron_h_mtp",
175+
"MLABlock": "mla",
176+
# gdn / kda / nsa / csa_hca / moe / mlp / engram / attention /
177+
# lightning_indexer are local closures in unified_superblock_v4.py
178+
# so their _SelfAttn / _MLP class names aren't unique. Callers
179+
# that need to walk such pre-built models should attach a
180+
# ``_v4_brick_kind`` attribute to the module instance, which we
181+
# honour below.
182+
}
183+
if hasattr(module, "_v4_brick_kind"):
184+
return getattr(module, "_v4_brick_kind")
185+
return KIND_BY_CLASS.get(cls_name)
186+
187+
188+
def from_mlx_model(
189+
model: nn.Module,
190+
*,
191+
attr_order: Sequence[str] | None = None,
192+
) -> BrickGraph:
193+
"""Walk ``model`` and extract BrickNodes from its direct children.
194+
195+
Only the IMMEDIATE children of ``model`` are considered (one level
196+
deep) — deeper recursion is the caller's job. This matches how
197+
UnifiedSuperblock composes bricks: as flat siblings under one parent.
198+
199+
Edges are inferred as a linear chain in the order children were
200+
declared (Python 3.7+ dict ordering). Override via ``attr_order``.
201+
"""
202+
nodes: list[BrickNode] = []
203+
children = (
204+
list(attr_order)
205+
if attr_order is not None
206+
else [name for name, _ in model.named_modules() if "." not in name and name]
207+
)
208+
used_names: set[str] = set()
209+
for child_name in children:
210+
child = getattr(model, child_name, None)
211+
if child is None or not isinstance(child, nn.Module):
212+
continue
213+
kind = _brick_kind_of(child)
214+
if kind is None:
215+
continue
216+
unique_name = child_name
217+
suffix = 0
218+
while unique_name in used_names:
219+
suffix += 1
220+
unique_name = f"{child_name}_{suffix}"
221+
used_names.add(unique_name)
222+
nodes.append(
223+
BrickNode(kind=kind, name=unique_name, params={}, module=child)
224+
)
225+
226+
edges = tuple(
227+
(nodes[i].name, nodes[i + 1].name) for i in range(len(nodes) - 1)
228+
)
229+
return BrickGraph(nodes=tuple(nodes), edges=edges)
230+
231+
232+
__all__ = [
233+
"BrickGraph",
234+
"BrickNode",
235+
"from_block_specs",
236+
"from_mlx_model",
237+
]

0 commit comments

Comments
 (0)