From b803a6857e48eb8aa5c4bbe2a8c9e4cf4eb30252 Mon Sep 17 00:00:00 2001 From: Lolo1222 Date: Thu, 27 Aug 2026 16:33:32 +0800 Subject: [PATCH] add passnet skills --- .../Passnet/skills/passnet-feedback/SKILL.md | 145 +++++ .../passnet-feedback/scripts/_common.py | 309 +++++++++++ .../passnet-feedback/scripts/analyze_graph.py | 224 ++++++++ .../passnet-feedback/scripts/check_pattern.py | 464 ++++++++++++++++ .../passnet-feedback/scripts/fetch_problem.py | 48 ++ .../scripts/parse_eval_log.py | 260 +++++++++ .../skills/passnet-orchestrate/SKILL.md | 524 ++++++++++++++++++ .../skills/passnet-pattern-fusion/SKILL.md | 333 +++++++++++ .../references/kernel-templates.md | 325 +++++++++++ .../Passnet/skills/passnet-skill/SKILL.md | 376 +++++++++++++ .../Passnet/skills/passnet-solve/SKILL.md | 288 ++++++++++ .../references/passbench-internals.md | 188 +++++++ .../skills/passnet-triton-opt/SKILL.md | 181 ++++++ 13 files changed, 3665 insertions(+) create mode 100644 skills/task-oriented/Passnet/skills/passnet-feedback/SKILL.md create mode 100644 skills/task-oriented/Passnet/skills/passnet-feedback/scripts/_common.py create mode 100644 skills/task-oriented/Passnet/skills/passnet-feedback/scripts/analyze_graph.py create mode 100644 skills/task-oriented/Passnet/skills/passnet-feedback/scripts/check_pattern.py create mode 100644 skills/task-oriented/Passnet/skills/passnet-feedback/scripts/fetch_problem.py create mode 100644 skills/task-oriented/Passnet/skills/passnet-feedback/scripts/parse_eval_log.py create mode 100644 skills/task-oriented/Passnet/skills/passnet-orchestrate/SKILL.md create mode 100644 skills/task-oriented/Passnet/skills/passnet-pattern-fusion/SKILL.md create mode 100644 skills/task-oriented/Passnet/skills/passnet-pattern-fusion/references/kernel-templates.md create mode 100644 skills/task-oriented/Passnet/skills/passnet-skill/SKILL.md create mode 100644 skills/task-oriented/Passnet/skills/passnet-solve/SKILL.md create mode 100644 skills/task-oriented/Passnet/skills/passnet-solve/references/passbench-internals.md create mode 100644 skills/task-oriented/Passnet/skills/passnet-triton-opt/SKILL.md diff --git a/skills/task-oriented/Passnet/skills/passnet-feedback/SKILL.md b/skills/task-oriented/Passnet/skills/passnet-feedback/SKILL.md new file mode 100644 index 000000000..f337705f0 --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-feedback/SKILL.md @@ -0,0 +1,145 @@ +--- +name: passnet-feedback +description: > + Fast, deterministic feedback for PassNet work: pre-flight pattern/match verification + WITHOUT burning a GPU evaluation, per-node bottleneck analysis of a sample's graphs, + and eval-log parsing into per-variant status + estimated score + failure classification. + Use BEFORE every GPU evaluation (check_pattern), at the START of a sample + (analyze_graph), and AFTER every evaluation (parse_eval_log). +--- + +GPU evaluations cost minutes and are rate-limited; these tools answer in seconds. The +iteration loop is: `analyze_graph` once → author passes → `check_pattern` until green → +GPU evaluate → `parse_eval_log` → fix the top issue → repeat. + +All scripts live in the `scripts/` directory next to this `SKILL.md`. In this repo, prefer +`SCRIPTS=/.claude/skills/passnet-feedback/scripts`. Invoke +`python3 $SCRIPTS/.py ...` from anywhere. They need: the PassNet repo importable +(auto-detected from the sample dir's `entry.sh` symlink, or set +`PYTHONPATH`/`PASSNET_ROOT`), torch, and for smoke/bench steps a GPU +(`CUDA_VISIBLE_DEVICES` respected; they degrade to CPU structure-checks without one). + +## 1. `analyze_graph.py` — where is the bottleneck, what is matchable + +```bash +python3 $SCRIPTS/analyze_graph.py --sample-dir [--bench] [--max-variants N] +# default --max-variants 3, dtype-spread +``` +- Captures the REAL dynamo graph of each variant (exactly what passes must match). +- Per node: op kind, target, written args/kwargs form, output shape/dtype, and callable + pattern **matchability** (`method`/`C-bound` = mirror exactly; `PY-sig positional` = + matchable; `PY-sig kwargs/partial` = normal callable pattern will normalize differently). + A high-value kwargs-form Python functional region may still be recoverable with an exact + manual FX `GraphModule` pattern; use passnet-pattern-fusion and confirm with + `check_pattern.py`. +- `--bench` (GPU): per-node eager µs via an instrumented FX interpreter + whole-forward + eager e2e estimate → tells you the absorbable time and a ROUGH speedup ceiling. The + ceiling assumes a ~70 µs fixed tax — real tax grows with FX node count (150–230 µs on + 13-node graphs); treat the printed ceiling as optimistic and recalibrate after your + first eval. +- Use the region table at the bottom as the starting plan: it lists maximal runs of + matchable nodes with their absorbed-µs totals. + +## 2. `check_pattern.py` — pre-flight a pass_dir in seconds (run before EVERY eval) + +```bash +python3 $SCRIPTS/check_pattern.py --sample-dir [--pass-dir ] [--smoke [--bench]] +# --max-variants N (default 8, dtype-spread; 0 = all). --bench implies --smoke. +# --smoke/--bench run on ONE variant per dtype; matching runs on all chosen variants. +``` +Replicates the harness pipeline faithfully and reports: +1. JSON manifest sanity (names ↔ files). +2. AST validation per pass (the real `validate_pass_source` when repo importable). +3. `replacement_func()` stability + **distinct-function count** (must be 1 when >1 pass — + otherwise passes WILL be silently dropped by the limit). +4. Pattern trace → real `SubgraphMatcher` against each variant's dynamo graph: + match count per pass per variant, single-output check, containment check, and on + mismatch a node-by-node nearest-miss diff (pattern form vs graph form). +5. Verdict per variant: would the run match (≥1 pass), or early-exit at 0.1. +6. `--smoke` (GPU): actually applies the passes (real `PassMgrBackend`) and runs one + poisoned warmup call + numeric comparison vs eager at the dtype's baseline tolerance — + catches Unauthorized-Operator, dtype mismatch, and gross numeric bugs pre-eval. +7. `--bench` (GPU): micro-benchmarks compiled-vs-eager e2e (200 calls) per variant for a + speedup preview (no dynamo guards, so real eval ≈ a few % worse). + +CPU-only output can validate structure and matchability triage, but it does not prove poisoned-wrapper legality, dtype behavior, numeric correctness, or real speed. GPU smoke/bench or a completed evaluation is still required for those claims. + +`--smoke` is a one-shot semantic check, not a full harness proof. Be extra cautious when a +replacement absorbs `inplace=True` nodes or other side-effectful behavior: a one-shot compare +can pass while repeated benchmark calls expose mutation/aliasing differences. For those +regions, either leave the in-place node outside the pattern or confirm repeated-call behavior +with a harness-style check/completed evaluation. + +Green output = safe to spend a GPU evaluation. Any red line tells you which skill to open: +match issues → passnet-pattern-fusion §7; poison/API issues → passnet-pattern-fusion §6; +numeric issues → passnet-triton-opt §3. + +## 3. `parse_eval_log.py` — turn an eval into decisions + +```bash +python3 $SCRIPTS/parse_eval_log.py +# also reads the "stdout"/"stderr" fields of a saved /evaluate JSON response +``` +Prints: +- Per-variant table: dtype | status | e2e/gpu speedup | eager/compiled medians | max_diff | + passes applied/failed. +- Failure classification with the fix pointer: + - localhost connection failure from a Codex managed sandbox → retry normal bounded curl once, then retry through the approved escalation path if available before diagnosing service downtime. + - empty, interrupted, HTTP000-style, status-corrupted, or non-JSON `/evaluate` response → no completed eval; leave score, speed, correctness, and pass-matched metrics null, retry service access/upload state if appropriate, and do not interpret any performance or correctness result. + - `no pass matched` + diagnostic lines → pattern form problem (which pass, which node). + - `AssertionError` in `_replace_pattern` → multi-output pattern (rule 1). + - `Unauthorized Operator (aten.xxx)` → illegal torch op in wrapper (poison). + - `Detected hacking behavior` → AST validation rejected a file (it was NOT loaded). + - `Loaded N passes` < listed → replacement_func limit dropped passes (shared dispatch!). + - dtype mismatch / accuracy with max_diff vs that dtype's baseline tolerance. + - replacement crash mentioning a returning node with no users → the pattern likely matched + userless/dead duplicate work; re-anchor the region on an observable value or use a proven + graph-rewrite approach instead of a normal output replacement. + - evaluator/environment fluctuation messages, timing-stability errors, or rerun requests + with otherwise clean matching and successful-variant correctness → classify separately + from numeric, no-match, unauthorized-operator, timeout, and OOM. If the pass pre-flights + cleanly and successful variants are correct, an unchanged rerun can be a valid use of + remaining evaluation budget. + - timeout/OOM hints. +- **Estimated sample score** using the real ES(t) weights, plus per-variant rectified + speedups — so you know the score impact of fixing each failing variant before re-running. + +For round or multi-worker aggregation, keep family labels separate: + +- `triage_family`: the static family assigned before implementation. +- `worker_confirmed_family`: the family the worker confirmed after reading the graph and + pre-flight results. +- `actual_winning_region_family`: the family of the best completed-evaluation state. + +Do not aggregate a win under the original triage family if the best completed state belongs +to a different region family. Round summaries should report, per family: number of completed +workers, score-above-eager count, correct-but-slow count, numeric-preflight-blocked count, +correctness-regressed-on-larger-region count, larger-region-attempted count, and evaluator +instability count. Keep evaluator/environment fluctuation separate from numeric, no-match, +unauthorized-operator, timeout, and OOM. + +## 4. Reading raw logs yourself (when needed) + +Key markers in `stdout`/`validation.log`: +`[PassMgrBackend] Loaded/Applied/failed to match/Diagnostic`, +`[Result] status: success|failed`, `[Speedup][e2e]/[gpu]`, +`[Performance][eager|compiled]` (medians), `[Correctness][max_diff|mean_diff|equal]`, +`[Datatype][eager|compiled]`, `debug-model-execution ` (crash), +`Has Any pass matched? [True|False]`, `aggregated_speedup=...`. +The service strips `Trial`/`[Profiling]`/`all_close` lines; everything above survives. + +## 5. Bottleneck decision guide (after analyze/eval) + +| signal | diagnosis | action | +|---|---|---| +| eager e2e < 150 µs, ≤2 matchable cheap nodes | overhead-bound, ceiling < 1 | floor pass, move on | +| big gap between eager e2e and Σ(node µs) | per-call fixed costs dominate | absorb more nodes per launch; nothing else helps | +| one node ≥ 60% of eager time, matchable | real kernel target | fuse it + its neighbors; tune via passnet-triton-opt | +| a `conv` dominates AND `stride == kernel_size`, `kernel_size > 1`, `padding == 0` (non-overlapping) | disjoint windows ⇒ exactly a dense patch matmul; cuDNN pays for a general im2col path | reformulate as a `tl.dot` patch-gather kernel + fused tail (kernel-templates §13), AFTER banking a tail floor. Don't dismiss as "cuDNN, leave alone" or "unmatchable" | +| a `conv`/`matmul` dominates in any OTHER form (1×1, depthwise/grouped, overlapping conv, general dense matmul) | vendor sweet spot | default: leave it in aten; fuse its cheap tail (bias/BN/act/residual). Rewriting usually loses — try only if parameters look off-regime and a completed eval beats the floor | +| one node ≥ 60%, callable-unmatchable kwargs form | try exact manual FX if the region is single-output and valuable; otherwise fuse what's left | expectation depends on whether the exact pattern pre-flight matches | +| compiled gpu ≈ e2e, both < 1 | stream gaps (launch-bound) | fewer launches/allocs | +| compiled gpu > 1, e2e < 1 | host overhead | drop autotune churn, simplify wrapper, fewer passes | +| some dtype variants fail only | numeric fidelity | passnet-triton-opt §3 recipes | +| timeout (600 s) | too many variants × compile/tune cost | remove autotune, single config, fewer passes | +| evaluator reports timing/environment fluctuation while successful variants match and pass correctness | evaluator instability, not proven kernel bug | rerun unchanged if budget remains; record separately from numeric/no-match/unauthorized | diff --git a/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/_common.py b/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/_common.py new file mode 100644 index 000000000..f9ac72aec --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/_common.py @@ -0,0 +1,309 @@ +"""Shared helpers for PassNet feedback scripts. + +Self-sufficient where possible; uses the real pass_bench implementations when the repo +is importable (preferred — identical semantics to the evaluator). +""" +import importlib.util +import inspect +import os +import re +import sys +from pathlib import Path + +import torch + + +# --------------------------------------------------------------------------- repo +def find_passnet_root(sample_dir=None): + """Locate the PassNet repo root (pass_bench importable).""" + env = os.environ.get("PASSNET_ROOT") + if env and (Path(env) / "pass_bench" / "__init__.py").exists(): + return Path(env) + if sample_dir is not None: + # entry.sh in a sample is a symlink into the repo + entry = Path(sample_dir) / "entry.sh" + if entry.exists(): + real = Path(os.path.realpath(str(entry))) + for p in real.parents: + if (p / "pass_bench" / "__init__.py").exists(): + return p + for p in Path(sample_dir).resolve().parents: + if (p / "pass_bench" / "__init__.py").exists(): + return p + try: + import pass_bench # noqa: F401 + return Path(pass_bench.__file__).resolve().parent.parent + except ImportError: + return None + + +def ensure_repo_on_path(sample_dir=None): + root = find_passnet_root(sample_dir) + if root is not None and str(root) not in sys.path: + sys.path.insert(0, str(root)) + return root + + +# --------------------------------------------------------------------------- graphs +def load_graph_list(sample_dir): + sample_dir = Path(sample_dir) + gl = sample_dir / "graph_list.txt" + if not gl.exists(): + raise FileNotFoundError(f"{gl} not found — is this a sample root?") + rels = [ln.strip() for ln in gl.read_text().splitlines() if ln.strip()] + return [(rel, (sample_dir / rel).resolve()) for rel in rels] + + +def variant_dtype(rel_path): + m = re.search(r"/(float16|float32|bfloat16|float64)/", rel_path) + return m.group(1) if m else "unknown" + + +def pick_variants(variants, max_variants): + """Pick up to max_variants, covering each dtype at least once first.""" + if max_variants is None or len(variants) <= max_variants: + return variants + by_dtype, picked, rest = {}, [], [] + for rel, d in variants: + dt = variant_dtype(rel) + if dt not in by_dtype: + by_dtype[dt] = (rel, d) + picked.append((rel, d)) + else: + rest.append((rel, d)) + for item in rest: + if len(picked) >= max_variants: + break + picked.append(item) + return picked[:max_variants] + + +def _modify_code_by_device(code, device): + try: + from pass_bench.torch.utils import modify_code_by_device + return modify_code_by_device(code, device) + except Exception: + # regex fallback: device(type='cuda'[, index=0]) and "cuda"/"cuda:0" strings + code = re.sub(r"device\(type='cuda'(?:,\s*index=\d+)?\)", f"device(type='{device}')", code) + code = re.sub(r"(['\"])cuda(?::\d+)?(['\"])", rf"\1{device}\2", code) + return code + + +def load_model(graph_dir, device): + """Load GraphModule class from model.py the way test_compiler does.""" + graph_dir = Path(graph_dir) + code = (graph_dir / "model.py").read_text() + dev_kind = "cuda" if str(device).startswith("cuda") else "cpu" + code = _modify_code_by_device(code, dev_kind) + spec = importlib.util.spec_from_loader(f"pn_model_{abs(hash(str(graph_dir)))}", loader=None) + module = importlib.util.module_from_spec(spec) + exec(compile(code, str(graph_dir / "model.py"), "exec"), module.__dict__) + model = module.GraphModule().to(torch.device(device)) + return model + + +# --------------------------------------------------------------------------- inputs +def parse_weight_meta(graph_dir): + """Parse weight_meta.py into {input_name: spec_dict}.""" + ns = {} + exec((Path(graph_dir) / "weight_meta.py").read_text(), ns) + specs = {} + for v in ns.values(): + if not isinstance(v, type) or not hasattr(v, "name"): + continue + specs[v.name] = { + "shape": list(getattr(v, "shape", [])), + "dtype": getattr(v, "dtype", "torch.float32"), + "device": getattr(v, "device", "cpu"), + "mean": getattr(v, "mean", 0.0), + "std": getattr(v, "std", 0.1), + "data": getattr(v, "data", None), + "min_val": getattr(v, "min_val", None), + "max_val": getattr(v, "max_val", None), + } + return specs + + +def replay_tensor(spec, device, force_dtype=None): + """Minimal replica of pass_bench.torch.utils.replay_tensor.""" + dtype = getattr(torch, spec["dtype"].replace("torch.", "")) + if force_dtype is not None and dtype.is_floating_point: + dtype = force_dtype + shape = spec["shape"] + if spec["data"] is not None: + t = torch.tensor(spec["data"], dtype=dtype).reshape(shape) + return t.to(device) + if dtype is torch.bool: + return (torch.randn(shape) > 0.5).to(dtype).to(device) + mean = spec["mean"] if spec["mean"] is not None else 0.0 + std = spec["std"] if spec["std"] is not None else 0.1 + if std == 0: + t = torch.full(shape, fill_value=mean, dtype=dtype) + else: + t = torch.randn(shape).to(dtype) * std * 0.2 + mean + if spec["min_val"] is not None: + t = torch.clamp(t, min=spec["min_val"]) + if spec["max_val"] is not None: + t = torch.clamp(t, max=spec["max_val"]) + if dtype.is_floating_point: + t = torch.where(torch.isfinite(t), t, torch.randn_like(t) * 0.01) + t = torch.clamp(t, min=-100.0, max=100.0) + return t.to(device) + + +def build_inputs(model, graph_dir, device, force_dtype=None, seed=123): + torch.manual_seed(seed) + specs = parse_weight_meta(graph_dir) + tensors = {k: replay_tensor(v, device, force_dtype) for k, v in specs.items()} + sig = inspect.signature(model.forward) + names = [n for n, p in sig.parameters.items() + if n != "self" and p.kind in (p.POSITIONAL_ONLY, p.POSITIONAL_OR_KEYWORD)] + missing = [n for n in names if n not in tensors] + if missing: + raise KeyError(f"weight_meta missing inputs {missing}; has {list(tensors)}") + return [tensors[n] for n in names] + + +# --------------------------------------------------------------------------- dynamo +def capture_dynamo_graph(model, inputs): + """Run torch.compile with a capture backend; return (gm, example_inputs).""" + import torch._dynamo as dynamo + dynamo.reset() + captured = {} + + def backend(gm, example_inputs): + captured["gm"] = gm + captured["inputs"] = example_inputs + return gm + + cm = torch.compile(model, backend=backend) + with torch.no_grad(): + torch.manual_seed(1024) + cm(*inputs) + if "gm" not in captured: + raise RuntimeError("dynamo did not capture a graph") + return captured["gm"], captured["inputs"] + + +def force_args_trace(fn): + """The harness's pattern tracer (normalizes call_function args).""" + try: + from pass_bench.torch.custom_replacement import force_args_symbolic_trace + return force_args_symbolic_trace(fn) + except ImportError: + pass + + class ForceArgsTracer(torch.fx.Tracer): + def create_node(self, kind, target, args, kwargs, name=None, type_expr=None): + if kind == "call_function" and callable(target): + try: + sig = inspect.signature(target) + bound = sig.bind(*args, **kwargs) + bound.apply_defaults() + return super().create_node(kind, target, tuple(bound.args), {}, name, type_expr) + except (ValueError, TypeError): + pass + return super().create_node(kind, target, args, kwargs, name, type_expr) + + tracer = ForceArgsTracer() + graph = tracer.trace(fn) + name = fn.__class__.__name__ if isinstance(fn, torch.nn.Module) else fn.__name__ + return torch.fx.GraphModule(tracer.root, graph, name) + + +# --------------------------------------------------------------------------- nodes +def target_name(node): + t = node.target + if node.op == "call_method": + return f".{t}" + if callable(t): + mod = getattr(t, "__module__", "") or "" + nm = getattr(t, "__name__", str(t)) + if mod.startswith("torch.nn.functional"): + return f"F.{nm}" + if mod == "torch" or mod.startswith("torch._C"): + return f"torch.{nm}" + if mod in ("_operator", "operator"): + return f"op.{nm}" + return f"{mod}.{nm}" if mod else nm + return str(t) + + +def fmt_arg(a, maxlen=28): + if isinstance(a, torch.fx.Node): + return f"%{a.name}" + s = repr(a) + return s if len(s) <= maxlen else s[: maxlen - 2] + ".." + + +def fmt_args(node): + parts = [fmt_arg(a) for a in node.args] + parts += [f"{k}={fmt_arg(v)}" for k, v in node.kwargs.items()] + return "(" + ", ".join(parts) + ")" + + +def classify_matchability(node): + """Return (tag, matchable: bool|None, note) for a dynamo target-graph node. + + Mirrors normal callable-pattern asymmetry: the pattern is force-args-normalized; the + target keeps the written form. A node is callable-matchable iff a pattern node can be + produced in the same form. Exact manual FX GraphModule patterns can recover some + kwargs-form Python functional nodes; analyze_graph keeps this classification + conservative so region suggestions are safe defaults. + """ + if node.op in ("placeholder", "output"): + return ("-", None, "") + if node.op == "get_attr": + return ("get_attr", True, "constant attr — pattern needs identical tensor type") + if node.op == "call_method": + return ("method", True, "mirror exactly (incl. kwargs)") + if node.op == "call_module": + return ("module", False, "call_module can't be expressed in a pattern fn") + target = node.target + try: + sig = inspect.signature(target) + except (ValueError, TypeError): + return ("C-bound", True, "mirror exactly (incl. kwargs)") + try: + bound = sig.bind(*node.args, **node.kwargs) + bound.apply_defaults() + normalized = tuple(bound.args) + except (ValueError, TypeError): + return ("PY-sig?", False, "args don't bind to signature — pattern can't reproduce") + if tuple(node.args) == normalized and not node.kwargs: + return ("PY-sig", True, "full-positional call — matchable") + return ("PY-sig", False, + "kwargs/omitted-defaults form — normal callable pattern normalizes differently; " + "manual FX may recover a valuable exact region") + + +RNG_TARGET_NAMES = {"torch.rand", "torch.randn", "torch.rand_like", "torch.randn_like", + "torch.randint", "torch.bernoulli", "torch.multinomial", + "torch.randperm", "torch.poisson", "torch.normal"} + + +def is_rng_node(node): + name = target_name(node) + if name in RNG_TARGET_NAMES: + return True + if name == "F.dropout" and node.op == "call_function": + # training flag: positional arg 2 or kwarg + training = None + if len(node.args) >= 3: + training = node.args[2] + training = node.kwargs.get("training", training) + return bool(training) is True + return False + + +def shape_of(node): + tm = node.meta.get("tensor_meta") + if tm is not None: + try: + return f"{list(tm.shape)}:{str(tm.dtype).replace('torch.', '')}" + except Exception: + pass + val = node.meta.get("example_value", None) + if isinstance(val, torch.Tensor): + return f"{list(val.shape)}:{str(val.dtype).replace('torch.', '')}" + return "" diff --git a/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/analyze_graph.py b/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/analyze_graph.py new file mode 100644 index 000000000..7aefff45a --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/analyze_graph.py @@ -0,0 +1,224 @@ +#!/usr/bin/env python3 +"""Analyze a PassNet sample: dynamo graph per variant, per-node matchability, +optional per-node eager timings (--bench), and fusion-region suggestions. + +Usage: + python3 analyze_graph.py --sample-dir [--bench] [--max-variants 3] + python3 analyze_graph.py --sample-dir ... +""" +import argparse +import sys +import time +from pathlib import Path + +sys.path.insert(0, str(Path(__file__).resolve().parent)) +from _common import ( # noqa: E402 + build_inputs, capture_dynamo_graph, classify_matchability, ensure_repo_on_path, + fmt_args, is_rng_node, load_graph_list, load_model, pick_variants, shape_of, + target_name, variant_dtype, +) + +import torch # noqa: E402 + + +class NodeTimer(torch.fx.Interpreter): + """Times each node with CUDA events over `reps` replays. + + Subtracts the measurement floor (sync + event overhead, measured on an empty loop) + from host times so Σ(node µs) approximates real absorbable time. + """ + + def __init__(self, gm, reps=30): + super().__init__(gm) + self.reps = reps + self.times = {} + self.floor_us = self._measure_floor() if torch.cuda.is_available() else 0.0 + + def _measure_floor(self): + start, end = torch.cuda.Event(True), torch.cuda.Event(True) + torch.cuda.synchronize() + t0 = time.perf_counter() + start.record() + for _ in range(self.reps): + pass + end.record() + torch.cuda.synchronize() + return (time.perf_counter() - t0) / self.reps * 1e6 + + def run_node(self, n): + if n.op in ("placeholder", "output", "get_attr") or not torch.cuda.is_available(): + return super().run_node(n) + # warmup once + result = super().run_node(n) + start, end = torch.cuda.Event(True), torch.cuda.Event(True) + torch.cuda.synchronize() + t0 = time.perf_counter() + start.record() + for _ in range(self.reps): + super().run_node(n) + end.record() + torch.cuda.synchronize() + host_us = (time.perf_counter() - t0) / self.reps * 1e6 + gpu_us = start.elapsed_time(end) * 1000 / self.reps + host_us = max(host_us - self.floor_us, gpu_us, 0.5) + self.times[n] = (host_us, gpu_us) + return result + + +def bench_eager_e2e(model, inputs, iters=200): + """Mirrors the harness model_call: manual_seed + forward + sync per trial.""" + with torch.no_grad(): + for _ in range(25): + model(*inputs) + torch.cuda.synchronize() + t0 = time.perf_counter() + for _ in range(iters): + torch.manual_seed(1024) + model(*inputs) + torch.cuda.synchronize() + return (time.perf_counter() - t0) / iters * 1e6 + + +def suggest_regions(gm, matchable_map, times): + """Connected components of matchable compute nodes (dataflow edges).""" + nodes = [n for n in gm.graph.nodes if matchable_map.get(n) and not is_rng_node(n)] + nodeset = set(nodes) + parent = {n: n for n in nodes} + + def find(x): + while parent[x] is not x: + parent[x] = parent[parent[x]] + x = parent[x] + return x + + def union(a, b): + ra, rb = find(a), find(b) + if ra is not rb: + parent[ra] = rb + + for n in nodes: + for a in n.all_input_nodes: + if a in nodeset: + union(n, a) + comps = {} + for n in nodes: + comps.setdefault(find(n), []).append(n) + regions = [] + for comp in comps.values(): + comp_set = set(comp) + outputs = [n for n in comp if any(u not in comp_set for u in n.users)] + total_us = sum(times.get(n, (0, 0))[0] for n in comp) if times else None + regions.append({"nodes": comp, "outputs": outputs, "us": total_us}) + regions.sort(key=lambda r: -(r["us"] or len(r["nodes"]))) + return regions + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--sample-dir", required=True) + ap.add_argument("--bench", action="store_true", help="per-node eager timings (GPU)") + ap.add_argument("--max-variants", type=int, default=3) + ap.add_argument("--device", default=None) + args = ap.parse_args() + + ensure_repo_on_path(args.sample_dir) + device = args.device or ("cuda" if torch.cuda.is_available() else "cpu") + force_dtype = None if device.startswith("cuda") else torch.float32 + if args.bench and not device.startswith("cuda"): + print("[warn] --bench needs CUDA; continuing without timings") + args.bench = False + + variants = load_graph_list(args.sample_dir) + print(f"sample: {args.sample_dir}") + print(f"variants: {len(variants)} " + f"({', '.join(sorted({variant_dtype(r) for r, _ in variants}))}); " + f"analyzing {min(args.max_variants, len(variants))} on {device}\n") + + for rel, gdir in pick_variants(variants, args.max_variants): + print("=" * 100) + print(f"variant: {rel}") + try: + model = load_model(gdir, device) + inputs = build_inputs(model, gdir, device, force_dtype) + gm, gm_inputs = capture_dynamo_graph(model, inputs) + except Exception as e: + print(f" [ERROR] could not capture graph: {type(e).__name__}: {e}") + continue + + # NB: gm placeholders follow dynamo's order — always feed gm with gm_inputs, + # never with the forward-signature-ordered `inputs`. + try: + from torch.fx.passes.shape_prop import ShapeProp + ShapeProp(gm).propagate(*gm_inputs) + except Exception: + pass + + times = {} + eager_e2e = None + if args.bench: + try: + timer = NodeTimer(gm) + with torch.no_grad(): + timer.run(*gm_inputs) + times = timer.times + eager_e2e = bench_eager_e2e(model, inputs) + except Exception as e: + print(f" [warn] bench failed: {type(e).__name__}: {e}") + + matchable_map = {} + print(f"{'#':>3} {'kind':13} {'target':26} {'written form':44} " + f"{'out shape':24} {'match?':7} {'µs':>8}") + idx = 0 + for n in gm.graph.nodes: + if n.op in ("placeholder", "output"): + continue + tag, ok, note = classify_matchability(n) + rng = is_rng_node(n) + matchable_map[n] = bool(ok) and not rng + us = f"{times[n][0]:8.1f}" if n in times else "" + flag = ("RNG!" if rng else ("yes" if ok else ("NO" if ok is False else "?"))) + print(f"{idx:>3} {n.op:13} {target_name(n):26} {fmt_args(n):44.44} " + f"{shape_of(n):24.24} {flag:7} {us}") + if ok is False or rng: + print(f" └─ {('RNG op — never include in a pattern' if rng else note)}") + idx += 1 + + if eager_e2e is not None: + node_sum = sum(t[0] for t in times.values()) + print(f"\n eager e2e ≈ {eager_e2e:.1f} µs/call (incl. the harness's per-call " + f"manual_seed+sync ≈60µs) Σ per-node µs ≈ {node_sum:.1f} " + f"(upper bound on absorbable; per-node timing adds sync overhead)") + + regions = suggest_regions(gm, matchable_map, times) + if regions: + print("\n candidate fusion regions (matchable connected components):") + for i, r in enumerate(regions): + names = " → ".join(target_name(n) for n in r["nodes"][:8]) + more = "" if len(r["nodes"]) <= 8 else f" (+{len(r['nodes']) - 8} more)" + us = f", ≈{r['us']:.0f} µs absorbable" if r["us"] else "" + multi = (" [!] region has ≥2 external outputs — split into one pass per output" + if len(r["outputs"]) > 1 else "") + print(f" R{i}: {len(r['nodes'])} nodes{us}: {names}{more}{multi}") + if eager_e2e is not None: + absorbable = sum(r["us"] or 0 for r in regions) + tax = 70.0 # guards+FX+wrapper+launch, calibrated on real evals + # assume a good fused kernel costs ~30% of the absorbed eager time + denom = max(eager_e2e - absorbable, 0.0) + tax + 0.30 * absorbable + ceiling = eager_e2e / denom + if ceiling < 1.05: + verdict = ("overhead-bound: ceiling < 1 — ship a floor pass, " + "don't chase >1") + elif ceiling < 1.2: + verdict = (f"marginal (rough ceiling ≈ {ceiling:.2f}x): expect ≈1.0; " + f"fuse the largest region, accept the result either way") + else: + verdict = f"fusible: rough ceiling ≈ {ceiling:.2f}x" + print(f"\n verdict: {verdict}") + else: + print("\n [!] no matchable compute nodes — only floor options are get_attr/" + "layout nodes; expect ≤1.0 and prioritize ANY match over speedup") + print() + + +if __name__ == "__main__": + main() diff --git a/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/check_pattern.py b/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/check_pattern.py new file mode 100644 index 000000000..60994c8ca --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/check_pattern.py @@ -0,0 +1,464 @@ +#!/usr/bin/env python3 +"""Pre-flight a PassNet pass_dir without spending a GPU evaluation. + +Replicates the harness pipeline: manifest -> AST validation -> pass loading -> +replacement_func stability/uniqueness -> pattern trace -> SubgraphMatcher against each +variant's real dynamo graph (+ nearest-miss diff) -> optional poisoned smoke run and +micro-benchmark. + +Usage: + python3 check_pattern.py --sample-dir [--pass-dir DIR] [--smoke] [--bench] + [--max-variants N] [--match-all] + +Exit code 0 = every analyzed variant has >=1 matching pass (and smoke passed, if used). +""" +import argparse +import inspect +import importlib.util +import json +import re +import sys +import time +import traceback +from pathlib import Path + +sys.path.insert(0, str(Path(__file__).resolve().parent)) +from _common import ( # noqa: E402 + build_inputs, capture_dynamo_graph, ensure_repo_on_path, fmt_args, force_args_trace, + load_graph_list, load_model, pick_variants, target_name, variant_dtype, +) + +import torch # noqa: E402 + +GREEN, RED, YEL, END = "\033[32m", "\033[31m", "\033[33m", "\033[0m" +def ok(s): return f"{GREEN}{s}{END}" +def bad(s): return f"{RED}{s}{END}" +def warn(s): return f"{YEL}{s}{END}" + +BASELINE_TOL = { # t=-5 (rtol, atol) + "torch.float32": (1.3e-6, 1e-5), + "torch.float16": (1e-3, 1e-5), + "torch.bfloat16": (1.6e-2, 1e-5), + "torch.float64": (1e-7, 1e-7), +} + + +def load_manifest(pass_dir): + mf = pass_dir / "sorted_output_pass_rule_names.json" + issues = [] + if not mf.exists(): + return None, [bad("sorted_output_pass_rule_names.json MISSING — nothing will load")] + try: + names = json.loads(mf.read_text()) + assert isinstance(names, list) and all(isinstance(x, str) for x in names) + except Exception as e: + return None, [bad(f"manifest unreadable: {e}")] + for n in names: + if not (pass_dir / f"{n}.py").exists(): + issues.append(bad(f"manifest lists '{n}' but {n}.py does not exist")) + listed = set(names) + for f in pass_dir.glob("*.py"): + if f.stem not in listed and not f.stem.startswith("_"): + issues.append(warn(f"{f.name} exists but is NOT in the manifest (won't load; " + f"prefix helpers with '_' to silence)")) + return names, issues + + +def ast_validate(pass_dir, names): + try: + from pass_bench.ast_util import validate_pass_source + except ImportError: + return {n: ["(pass_bench not importable — AST validation skipped)"] for n in names}, False + out = {} + for n in names: + p = pass_dir / f"{n}.py" + if not p.exists(): + continue + try: + out[n] = validate_pass_source(p.read_text()) + except SyntaxError as e: + out[n] = [f"SyntaxError: {e}"] + return out, True + + +def load_pass_modules(pass_dir, names): + """Load pass modules like the harness (sample root on sys.path for pass_dir.* imports).""" + sample_root = str(pass_dir.parent) + if sample_root not in sys.path: + sys.path.insert(0, sample_root) + if pass_dir.name != "pass_dir": + text = "".join((pass_dir / f"{n}.py").read_text() + for n in names if (pass_dir / f"{n}.py").exists()) + if "pass_dir." in text or "from pass_dir" in text: + print(warn(f"pass dir is named '{pass_dir.name}' but passes import " + f"'pass_dir.*' — name the directory 'pass_dir' (e.g. " + f"/tmp/ws/pass_dir) or imports will fail here AND in the harness")) + mods, errors = {}, {} + for n in names: + p = pass_dir / f"{n}.py" + if not p.exists(): + continue + try: + spec = importlib.util.spec_from_file_location(n, p) + m = importlib.util.module_from_spec(spec) + m.__file__ = str(p) + spec.loader.exec_module(m) + for fn in ("pattern", "replacement_args", "replacement_func"): + if not callable(getattr(m, fn, None)): + raise AttributeError(f"missing function {fn}()") + mods[n] = m + except Exception as e: + errors[n] = f"{type(e).__name__}: {e}" + return mods, errors + + +def check_replacement_funcs(mods): + msgs, funcs = [], {} + for n, m in mods.items(): + try: + f1, f2 = m.replacement_func(), m.replacement_func() + if f1 is not f2: + msgs.append(bad(f"{n}: replacement_func() UNSTABLE (returns new object per " + f"call) — harness raises; return a module-level function")) + if not callable(f1): + msgs.append(bad(f"{n}: replacement_func() returned non-callable {f1!r}")) + funcs[n] = f1 + except Exception as e: + msgs.append(bad(f"{n}: replacement_func() raised {type(e).__name__}: {e}")) + distinct = {id(f) for f in funcs.values()} + if len(funcs) > 1 and len(distinct) > 1: + msgs.append(bad( + f"{len(funcs)} passes but {len(distinct)} distinct replacement functions — " + f"output_pass_replacement_func_limit=1 will RANDOMLY DROP all but one pass. " + f"All passes must `from pass_dir._shared_kernels import dispatch_wrapper` " + f"and return that same object (see passnet-pattern-fusion §4).")) + return msgs, funcs + + +def single_output_check(pattern_gm): + out_node = next(n for n in pattern_gm.graph.nodes if n.op == "output") + arg = out_node.args[0] + if isinstance(arg, (tuple, list)): + n_out = len(arg) + else: + n_out = 1 + return n_out, len(out_node.all_input_nodes) + + +def nearest_miss(pattern_gm, target_gm, max_lines=8): + """Per pattern compute node, show target nodes with the same target (form diff).""" + lines = [] + pat_nodes = [n for n in pattern_gm.graph.nodes + if n.op not in ("placeholder", "output")] + tgt_nodes = [n for n in target_gm.graph.nodes + if n.op not in ("placeholder", "output")] + for pn in pat_nodes: + same = [tn for tn in tgt_nodes if tn.op == pn.op and tn.target == pn.target] + if not same: + kind = {"call_function": "function", "call_method": "method"}.get(pn.op, pn.op) + others = [tn for tn in tgt_nodes + if target_name(tn).split(".")[-1] == target_name(pn).split(".")[-1]] + hint = (f" — graph has same-named node as {others[0].op} " + f"{target_name(others[0])}{fmt_args(others[0])} (form mismatch!)" + if others else " — no node with this target in the graph at all") + lines.append(f" pattern {kind} {target_name(pn)}{fmt_args(pn)}: " + f"NO target-node{hint}") + else: + pa = fmt_args(pn) + forms = {fmt_args(tn) for tn in same} + if pa not in forms: + lines.append(f" pattern {target_name(pn)}{pa} vs graph " + f"{' | '.join(sorted(forms)[:3])} ← arg/kwarg/literal diff") + if len(lines) >= max_lines: + break + if not lines: + lines.append(" every pattern node has a same-form twin — failure is structural: " + "check dataflow edges, containment (an intermediate is consumed " + "outside the pattern), or overlapping matches") + return lines + + +def run_matcher(pattern_gm, target_gm): + from torch.fx.passes.utils.matcher_utils import SubgraphMatcher + m = SubgraphMatcher(pattern_gm.graph, match_output=False, match_placeholder=False, + remove_overlapping_matches=True, ignore_literals=False) + return m.match(target_gm.graph) + + +def pattern_graph_for(pattern): + """Mirror the harness: GraphModule/Graph patterns bypass ForceArgsTracer.""" + if isinstance(pattern, torch.fx.GraphModule): + return pattern + if isinstance(pattern, torch.fx.Graph): + return torch.fx.GraphModule({}, pattern, "ManualPattern") + return force_args_trace(pattern) + + +def pattern_signature_params(pattern, pattern_gm): + """Return parameter names used by PassMgrBackend for replacement wrapper generation.""" + if isinstance(pattern, torch.fx.Graph): + return [], [n.name for n in pattern_gm.graph.nodes if n.op == "placeholder"] + sig = inspect.signature(pattern) + params = list(sig.parameters) + placeholders = [n.name for n in pattern_gm.graph.nodes if n.op == "placeholder"] + if isinstance(pattern, (torch.fx.GraphModule, torch.fx.Graph)) and params != placeholders: + return params, placeholders + return params, [] + + +def smoke_variant(rel, gdir, device, pass_dir): + """Apply passes with the real PassMgrBackend; one poisoned call; numeric compare.""" + import pass_bench.torch.backend.pass_mgr_backend as pmb + from pass_bench.torch.backend.pass_mgr_backend import PassMgrBackend + from pass_bench.torch.override_dispatch_flag import global_override_dispatch + # The harness runs each variant in a fresh subprocess; mirror that isolation for the + # process-global replacement-function registry (each load creates new func objects). + pmb.g_replacement_func = None + model = load_model(gdir, device) + inputs = build_inputs(model, gdir, device) + backend = PassMgrBackend({ + "input_pass_rule_dir": str(Path(gdir).parents[0] / "__none__"), + "output_pass_rule_dir": str(pass_dir), + "output_pass_pattern_limit": 100, + "output_pass_replacement_func_limit": 1, + "pass_match_result_file_path": None, + }) + import torch._dynamo as dynamo + dynamo.reset() + cm = backend(model) + with torch.no_grad(): + torch.manual_seed(1024) + with global_override_dispatch(True): # anti-cheat warmup parity + compiled_out = cm(*inputs) + with global_override_dispatch(False): + torch.manual_seed(1024) + compiled_out = cm(*inputs) + torch.manual_seed(1024) + eager_out = model(*inputs) + if not isinstance(compiled_out, tuple): + compiled_out = (compiled_out,) + if not isinstance(eager_out, tuple): + eager_out = (eager_out,) + msgs, all_ok = [], True + for i, (e, c) in enumerate(zip(eager_out, compiled_out)): + if not isinstance(e, torch.Tensor): + continue + if e.dtype != c.dtype: + msgs.append(bad(f"out[{i}] dtype mismatch: eager {e.dtype} vs compiled {c.dtype}")) + all_ok = False + continue + rtol, atol = BASELINE_TOL.get(str(e.dtype), (1e-5, 1e-6)) + close = torch.allclose(e.float(), c.float(), rtol=rtol, atol=atol) + md = (e.float() - c.float()).abs().max().item() + line = f"out[{i}] {str(e.dtype).replace('torch.','')} max_diff={md:.3e} baseline({rtol:g},{atol:g}) → {'OK' if close else 'FAIL'}" + msgs.append(ok(line) if close else bad(line)) + all_ok &= close + return all_ok, msgs, (model, cm, inputs) + + +def bench_pair(model, cm, inputs, iters=200): + """Mirrors the harness model_call: manual_seed + forward + sync per trial.""" + def run(fn): + with torch.no_grad(): + for _ in range(25): + fn(*inputs) + torch.cuda.synchronize() + t0 = time.perf_counter() + for _ in range(iters): + torch.manual_seed(1024) + fn(*inputs) + torch.cuda.synchronize() + return (time.perf_counter() - t0) / iters * 1e6 + e, c = run(model), run(cm) + return e, c, e / c + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--sample-dir", required=True) + ap.add_argument("--pass-dir", default=None) + ap.add_argument("--max-variants", type=int, default=8, + help="match-check this many variants, dtype-spread (default 8; " + "0 = all — beware samples with 100+ variants)") + ap.add_argument("--smoke", action="store_true", help="apply passes + numeric smoke (GPU)") + ap.add_argument("--bench", action="store_true", + help="micro-benchmark eager vs compiled (implies --smoke; GPU)") + ap.add_argument("--device", default=None) + args = ap.parse_args() + if args.bench: + args.smoke = True + + sample_dir = Path(args.sample_dir).resolve() + pass_dir = Path(args.pass_dir).resolve() if args.pass_dir else sample_dir / "pass_dir" + repo = ensure_repo_on_path(sample_dir) + device = args.device or ("cuda" if torch.cuda.is_available() else "cpu") + force_dtype = None if device.startswith("cuda") else torch.float32 + failures = 0 + + print(f"sample: {sample_dir}\npass_dir: {pass_dir}\nrepo: {repo}\ndevice: {device}\n") + + # 1. manifest + names, issues = load_manifest(pass_dir) + for line in issues: + print(line) + if "MISSING" in line or "unreadable" in line or "does not exist" in line: + failures += 1 + if names is None: + sys.exit(1) + print(f"manifest: {names}") + + # 2. AST validation + ast_results, real_ast = ast_validate(pass_dir, names) + for n, viols in ast_results.items(): + if viols and real_ast: + failures += 1 + print(bad(f"AST validation FAILED for {n} (harness will SKIP this pass):")) + for v in viols: + print(f" - {v}") + elif viols: + print(warn(f"{n}: {viols[0]}")) + else: + print(ok(f"AST validation OK for {n}")) + + # 3. load + replacement funcs + mods, load_errors = load_pass_modules(pass_dir, names) + for n, err in load_errors.items(): + failures += 1 + print(bad(f"IMPORT FAILED for {n}: {err}")) + rf_msgs, _funcs = check_replacement_funcs(mods) + for m in rf_msgs: + print(m) + failures += 1 + + # 4. pattern traces + patterns = {} + for n, m in mods.items(): + try: + pg = pattern_graph_for(m.pattern) + n_out, _ = single_output_check(pg) + if n_out != 1: + failures += 1 + print(bad(f"{n}: pattern returns {n_out} values — harness CRASHES on " + f"multi-output patterns; split into {n_out} passes")) + continue + pat_args, placeholder_mismatch = pattern_signature_params(m.pattern, pg) + rep_args = list(inspect.signature(m.replacement_args).parameters) + n_pat_args = len(pat_args) + n_rep_args = len(rep_args) + if isinstance(m.pattern, torch.fx.Graph): + failures += 1 + print(bad(f"{n}: bare torch.fx.Graph patterns do not expose a callable " + f"signature for the backend wrapper — wrap it in " + f"torch.fx.GraphModule and set pattern.__signature__")) + elif placeholder_mismatch: + failures += 1 + print(bad(f"{n}: GraphModule pattern signature {pat_args} does not match " + f"graph placeholders {placeholder_mismatch} — set " + f"pattern.__signature__ to the placeholder order")) + if n_pat_args != n_rep_args: + failures += 1 + print(bad(f"{n}: pattern takes {n_pat_args} args but replacement_args " + f"takes {n_rep_args} — must be identical")) + patterns[n] = pg + print(ok(f"{n}: pattern traced, single output, " + f"{sum(1 for x in pg.graph.nodes if x.op not in ('placeholder', 'output'))} nodes")) + except Exception as e: + failures += 1 + print(bad(f"{n}: pattern trace FAILED: {type(e).__name__}: {e}")) + + # 5. match per variant + variants = load_graph_list(sample_dir) + chosen = pick_variants(variants, args.max_variants or None) + if len(chosen) < len(variants): + skipped = [r for r, _ in variants if r not in {c for c, _ in chosen}] + print(f"\n[note] match-checking {len(chosen)}/{len(variants)} variants " + f"(--max-variants {args.max_variants}); dtype coverage guaranteed, " + f"skipped same-dtype extras:") + for r in skipped[:6]: + print(f" - {r}") + if len(skipped) > 6: + print(f" ... +{len(skipped) - 6} more (pass --max-variants 0 for all)") + print(f"\nmatching {len(patterns)} pass(es) against {len(chosen)}/{len(variants)} variants:") + unmatched_variants = [] + target_cache = {} + for rel, gdir in chosen: + try: + model = load_model(gdir, device) + inputs = build_inputs(model, gdir, device, force_dtype) + gm, _ = capture_dynamo_graph(model, inputs) + target_cache[rel] = (gm, gdir) + except Exception as e: + print(bad(f" {rel}: dynamo capture failed: {type(e).__name__}: {e}")) + traceback.print_exc(limit=2) + unmatched_variants.append(rel) + continue + any_match = False + per = [] + for n, pg in patterns.items(): + try: + ms = run_matcher(pg, gm) + except Exception as e: + per.append(f"{n}: matcher error {e}") + continue + per.append(f"{n}: {len(ms)}") + any_match |= bool(ms) + status = ok("MATCH") if any_match else bad("NO MATCH → variant would score 0.1") + seed_m = re.search(r"/(float\d+|bfloat16)/(\w+)/", rel) + vtag = f"{seed_m.group(1)}/{seed_m.group(2)}" if seed_m else variant_dtype(rel) + print(f" [{vtag:12}] {rel.split('/')[-1][:44]:44} {status} ({', '.join(per)})") + if not any_match: + unmatched_variants.append(rel) + for n, pg in patterns.items(): + print(f" nearest-miss for {n}:") + for line in nearest_miss(pg, gm): + print(line) + + if unmatched_variants: + failures += len(unmatched_variants) + + # 6. smoke + if args.smoke and patterns and not unmatched_variants: + if not device.startswith("cuda"): + print(warn("\n--smoke needs CUDA; skipped")) + elif repo is None: + print(warn("\n--smoke needs pass_bench importable; skipped")) + else: + print("\nsmoke (real PassMgrBackend, poisoned warmup, numeric compare):") + seen_dtypes = set() + for rel, gdir in chosen: + dt = variant_dtype(rel) + if dt in seen_dtypes: + continue + seen_dtypes.add(dt) + try: + good, msgs, trio = smoke_variant(rel, gdir, device, pass_dir) + for m in msgs: + print(f" [{dt:8}] {m}") + if not good: + failures += 1 + elif args.bench: + e, c, s = bench_pair(*trio) + line = (f" [{dt:8}] micro-bench eager {e:.1f}µs vs compiled {c:.1f}µs " + f"→ ~{s:.2f}x (full eval adds guard overhead)") + print(ok(line) if s >= 1 else warn(line)) + except Exception as e: + failures += 1 + print(bad(f" [{dt:8}] smoke CRASHED: {type(e).__name__}: {e}")) + tb = traceback.format_exc() + if "Unauthorized Operator" in tb: + print(bad(" → poison dispatch: an illegal aten op runs inside " + "your wrapper (only empty/zeros/ones/full/as_tensor/.to " + "+ metadata are allowed; do ALL math in Triton)")) + else: + print(" " + tb.strip().splitlines()[-1]) + + print() + if failures == 0: + print(ok(f"PRE-FLIGHT GREEN — safe to spend a GPU evaluation.")) + else: + print(bad(f"PRE-FLIGHT: {failures} blocking issue(s) — fix before evaluating.")) + sys.exit(0 if failures == 0 else 1) + + +if __name__ == "__main__": + main() diff --git a/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/fetch_problem.py b/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/fetch_problem.py new file mode 100644 index 000000000..8b23d2b3d --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/fetch_problem.py @@ -0,0 +1,48 @@ +#!/usr/bin/env python3 +"""Materialize a service-mode PassNet problem into a local directory so that +analyze_graph.py / check_pattern.py can run on it. + +Usage: + python3 fetch_problem.py --svc http://127.0.0.1:8765 --sample --out /tmp/ws + python3 fetch_problem.py --svc ... --out /tmp/ws # API-server mode (no --sample) + +Creates: /graph_list.txt, /graphs/.../{model.py,weight_meta.py}, /pass_dir/ +Then author passes in /pass_dir, pre-flight with check_pattern.py --sample-dir , +and POST the files to the service when green. +""" +import argparse +import json +import urllib.parse +import urllib.request +from pathlib import Path + + +def get(url): + with urllib.request.urlopen(url, timeout=30) as r: + return json.loads(r.read().decode()) + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--svc", required=True) + ap.add_argument("--sample", default=None, help="sample_path (service mode)") + ap.add_argument("--out", required=True) + args = ap.parse_args() + + q = f"?sample_path={urllib.parse.quote(args.sample, safe='')}" if args.sample else "" + prob = get(f"{args.svc}/problem{q}") + + out = Path(args.out) + (out / "pass_dir").mkdir(parents=True, exist_ok=True) + (out / "graph_list.txt").write_text("\n".join(prob["graph_list"]) + "\n") + for g in prob["graphs"]: + gdir = out / g["name"] + gdir.mkdir(parents=True, exist_ok=True) + (gdir / "model.py").write_text(g.get("model_code") or "") + (gdir / "weight_meta.py").write_text(g.get("weight_meta") or "") + print(f"materialized {len(prob['graphs'])} graphs under {out}") + print(f"next: python3 analyze_graph.py --sample-dir {out}") + + +if __name__ == "__main__": + main() diff --git a/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/parse_eval_log.py b/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/parse_eval_log.py new file mode 100644 index 000000000..9e5eb2a95 --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-feedback/scripts/parse_eval_log.py @@ -0,0 +1,260 @@ +#!/usr/bin/env python3 +"""Parse a PassNet evaluation output into per-variant results, failure classification, +and an estimated sample score. + +Accepts: a raw validation.log, a pass_evaluator stdout capture, or a saved JSON response +from POST /evaluate (uses its "stdout"+"stderr" fields and reports its "score"). + +Usage: python3 parse_eval_log.py [--sample-dir DIR] +""" +import argparse +import json +import math +import re +import sys +from pathlib import Path + +sys.path.insert(0, str(Path(__file__).resolve().parent)) +from _common import ensure_repo_on_path # noqa: E402 + +BASELINE_TOL = { + "float32": (1.3e-6, 1e-5), + "float16": (1e-3, 1e-5), + "bfloat16": (1.6e-2, 1e-5), + "float64": (1e-7, 1e-7), +} +# weights from aggregate_es_scores.get_weights() +WEIGHTS = {-10: 0.001, -9: 0.001, -8: 0.001, -7: 0.001, -6: 0.001, -5: 1.0, -4: 1.0, + -3: 1.0, -2: 0.8, -1: 0.64, 0: 0.512, 1: 0.4096, 2: 0.32768, 3: 0.262144, + 4: 0.001} +WSUM = sum(WEIGHTS.values()) + + +def read_text(path): + raw = Path(path).read_text(errors="ignore") + score = None + stripped = raw.lstrip() + if stripped.startswith("{"): + try: + j = json.loads(raw) + txt = (j.get("stdout", "") or "") + "\n" + (j.get("stderr", "") or "") + score = j.get("score") + if j.get("error"): + txt += f"\n[service-error] {j['error']}" + return txt, score, j + except json.JSONDecodeError: + pass + return raw, score, None + + +def split_per_variant(text): + blocks, cur = [], [] + for line in text.splitlines(): + if "[Processing]" in line: + if cur: + blocks.append(cur) + cur = [line] + elif cur: + cur.append(line) + if cur: + blocks.append(cur) + return blocks + + +F_SPEED = re.compile(r"\[Speedup\]\[(\w+)\]: ([\d.eE+-]+)") +F_PERF = re.compile(r"\[Performance\]\[(\w+)\]: (\{.*\})") +F_MAXD = re.compile(r"\[Correctness\]\[max_diff\]: (.+)") +F_DTYPE = re.compile(r"\[Datatype\]\[(\w+)\]: (.+)") +F_RESULT = re.compile(r"\[Result\] status: (\w+)") +F_APPLIED = re.compile(r"Applied (\d+) replacements with (\w+)") +F_FAILMATCH = re.compile(r"Pass (\w+) failed to match") +F_ALLCLOSE = re.compile(r"\[Correctness\]\[all_close_atol_([\d.E+-]+)_rtol_([\d.E+-]+)\]: (.+)") + + +def parse_block(lines): + d = {"path": lines[0].split()[-1], "speed": {}, "applied": [], "failed_match": [], + "max_diff": None, "dtype": None, "status": None, "errors": [], "allclose": {}, + "eager_med": None, "compiled_med": None, "diag": []} + for ln in lines: + if m := F_SPEED.search(ln): + d["speed"][m.group(1)] = float(m.group(2)) + elif m := F_PERF.search(ln): + try: + d[f"{m.group(1)}_med"] = json.loads(m.group(2))["e2e"]["median"] + except Exception: + pass + elif m := F_MAXD.search(ln): + d["max_diff"] = [float(x) for x in m.group(1).split()] + elif m := F_DTYPE.search(ln): + if m.group(1) == "eager": + d["dtype"] = m.group(2).split()[0] + elif m := F_RESULT.search(ln): + d["status"] = m.group(1) + elif m := F_APPLIED.search(ln): + d["applied"].append((m.group(2), int(m.group(1)))) + elif m := F_FAILMATCH.search(ln): + d["failed_match"].append(m.group(1)) + elif m := F_ALLCLOSE.search(ln): + vals = m.group(3).split() + d["allclose"][(float(m.group(1)), float(m.group(2)))] = all(v == "1" for v in vals) + elif "Unauthorized Operator" in ln: + d["errors"].append("poison: illegal aten op in wrapper") + elif "Detected hacking behavior" in ln: + d["errors"].append("AST validation rejected a pass file (not loaded)") + elif "No passes modified the graph" in ln: + d["errors"].append("no pass matched this variant") + elif "AssertionError" in ln and "custom_replacement" in "".join(lines): + d["errors"].append("harness assert — likely multi-output pattern") + elif "CUDA out of memory" in ln: + d["errors"].append("OOM") + elif "illegal memory access" in ln.lower(): + d["errors"].append("CUDA illegal memory access (kernel indexing bug)") + elif "debug-model-execution" in ln: + d["errors"].append(f"crash: {ln.split()[1]}") + elif "Diagnostic for" in ln or re.match(r"\s+- MatchFailure", ln): + d["diag"].append(ln.strip()) + elif m := re.search(r"Loaded (\d+) passes: (\[.*\])", ln): + d["loaded"] = (int(m.group(1)), m.group(2)) + return d + + +def baseline_correct(d): + """correct at t=-5 from allclose keys (unfiltered logs) or heuristics.""" + if d["status"] != "success": + return False, "exec-failed" + dt = (d["dtype"] or "").replace("torch.", "") + rtol, atol = BASELINE_TOL.get(dt, (1e-5, 1e-6)) + for (a, r), is_ok in d["allclose"].items(): + if math.isclose(a, atol, rel_tol=0.01) and math.isclose(r, rtol, rel_tol=0.01): + return is_ok, "exact" + if d["max_diff"] is not None: # filtered log: heuristic + md = max(d["max_diff"]) + if md == 0.0: + return True, "max_diff=0" + return md <= atol, f"heuristic(max_diff={md:.2e} vs atol={atol:g}; verify)" + return True, "assumed(status=success)" + + +def estimate_score(blocks): + """ES(t)-weighted estimate. Failure→0.1; accuracy-failure tolerated at t>=1.""" + per_t = {} + for t in WEIGHTS: + vals = [] + for d in blocks: + okv, _ = baseline_correct(d) + s = d["speed"].get("e2e") + if d["status"] == "success" and okv and s: + vals.append(s) # approximation: correct@-5 ⇒ correct for t≥-5 + if t < -5: + # strict levels usually fail unless bit-exact + md = max(d["max_diff"]) if d["max_diff"] else 1.0 + vals[-1] = s if md == 0.0 else 0.1 + else: + tolerated = (t >= 1 and d["status"] == "success") or t >= 3 + vals.append(1.0 if tolerated else 0.1) + if vals: + per_t[t] = math.exp(sum(math.log(max(v, 1e-10)) for v in vals) / len(vals)) + if not per_t: + return None + logscore = sum(WEIGHTS[t] * math.log10(per_t[t]) for t in per_t) / WSUM + return 10 ** logscore + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("logfile") + ap.add_argument("--sample-dir", default=None, help="for repo autodetection") + args = ap.parse_args() + ensure_repo_on_path(args.sample_dir) + + text, svc_score, svc_json = read_text(args.logfile) + blocks = [parse_block(b) for b in split_per_variant(text)] + if not blocks: + print("no [Processing] blocks found — the run died before evaluating any variant.") + if "No replacement functions available after filtering 0 rules" in text: + print("cause: pass_dir effectively EMPTY — no pass loaded (missing manifest, " + "AST-rejected files, or no .py files). Sample scores 0.1. " + "Run check_pattern.py to see why each pass was dropped.") + elif "failed to match" in text or "No passes modified the graph" in text: + print("cause: pass(es) loaded but matched NOTHING. Diagnostic lines:") + for ln in text.splitlines(): + if "failed to match" in ln or "Diagnostic" in ln or "MatchFailure" in ln: + print(f" {ln.strip()}") + print("→ run check_pattern.py for nearest-miss diffs " + "(passnet-pattern-fusion §7).") + elif "Detected hacking behavior" in text: + print("cause: AST validation rejected your pass file(s) — see " + "passnet-pattern-fusion §2 rule 9.") + elif "ModuleNotFoundError" in text or "ImportError" in text: + print("cause: a pass file failed to import:") + for ln in text.splitlines(): + if "Error" in ln: + print(f" {ln.strip()}") + break + elif "timed out" in text.lower(): + print("cause: evaluation timeout (600 s) — too many variants × compile cost; " + "remove autotune, reduce pass count.") + if svc_json is not None: + print(f"service fields: returncode={svc_json.get('returncode')} " + f"pass_matched={svc_json.get('pass_matched')} score={svc_score} " + f"error={svc_json.get('error')!r}") + sys.exit(1) + + print(f"{'dtype':9} {'status':9} {'e2e':>6} {'gpu':>6} {'eager':>8} {'comp':>8} " + f"{'max_diff':>10} {'corr@-5':14} notes") + n_corr = 0 + speeds = [] + for d in blocks: + okv, how = baseline_correct(d) + n_corr += okv + dt = (d["dtype"] or "?").replace("torch.", "") + s_e2e = d["speed"].get("e2e") + if okv and s_e2e: + speeds.append(s_e2e) + notes = [] + if d["failed_match"]: + notes.append(f"no-match:{','.join(d['failed_match'])}") + if d["applied"]: + notes.append("applied:" + ",".join(f"{n}×{c}" for n, c in d["applied"])) + notes += d["errors"][:2] + print(f"{dt:9} {(d['status'] or 'crash'):9} " + f"{s_e2e or float('nan'):6.3f} {d['speed'].get('gpu', float('nan')):6.3f} " + f"{(d.get('eager_med') or float('nan')) * 1000:8.1f} " + f"{(d.get('compiled_med') or float('nan')) * 1000:8.1f} " + f"{(max(d['max_diff']) if d['max_diff'] else float('nan')):10.2e} " + f"{('OK(' + how + ')' if okv else 'FAIL(' + how + ')'):14.14} " + f"{'; '.join(notes)}") + for ln in d["diag"][:4]: + print(f" diag: {ln}") + + print(f"\nvariants: {len(blocks)} correct@baseline: {n_corr}/{len(blocks)}") + if speeds: + gm = math.exp(sum(math.log(s) for s in speeds) / len(speeds)) + print(f"gmean e2e speedup over correct variants: {gm:.3f} " + f"fast_1: {sum(s >= 1 for s in speeds)}/{len(speeds)}") + m = re.search(r"aggregated_speedup=([\d.eE+-]+)", text) + if m: + print(f"authoritative aggregated score (from log): {float(m.group(1)):.4f}") + elif svc_score is not None: + sc = svc_score.get("score") if isinstance(svc_score, dict) else svc_score + print(f"authoritative score (service): {sc}") + est = estimate_score(blocks) + if est is not None: + print(f"estimated sample score (ES-weighted): {est:.4f}") + # impact preview: each failing variant fixed to speedup 1.0 + fails = [i for i, d in enumerate(blocks) + if not baseline_correct(d)[0] or not d["speed"].get("e2e")] + if fails: + import copy + fixed_blocks = copy.deepcopy(blocks) + for i in fails: + fixed_blocks[i]["status"] = "success" + fixed_blocks[i]["speed"]["e2e"] = 1.0 + fixed_blocks[i]["max_diff"] = [0.0] + est2 = estimate_score(fixed_blocks) + print(f"if the {len(fails)} failing variant(s) were fixed to 1.0x: ≈{est2:.4f} " + f"(+{est2 - est:.4f}) ← prioritize fixing failures over tuning winners") + + +if __name__ == "__main__": + main() diff --git a/skills/task-oriented/Passnet/skills/passnet-orchestrate/SKILL.md b/skills/task-oriented/Passnet/skills/passnet-orchestrate/SKILL.md new file mode 100644 index 000000000..16fc7decf --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-orchestrate/SKILL.md @@ -0,0 +1,524 @@ +--- +name: passnet-orchestrate +description: > + Round/batch planning and decision support for PassNet graph optimization. Use for + multi-sample triage, worker allocation, eval-budget policy, and post-eval keep/revert + decisions. For solving one concrete sample end to end, use passnet-solve as the entry + skill; consult this skill only when you need broader planning or a specific decision gate. +--- + +This skill is the **planning and decision** layer for PassNet. It is not the single-sample +entry point. For one concrete sample, start with **passnet-solve**; it owns the end-to-end +loop and calls passnet-pattern-fusion, passnet-triton-opt, and passnet-feedback at the right +time. Use this skill when planning a multi-sample round, choosing which sample families to +dispatch, auditing whether a line of attack is overhead-bound, or deciding whether to keep, +revert, or stop after completed evaluations. + +For context, a **sample** = several **graph variants** of the same subgraph (different dtypes — +float32 / float16 / bfloat16 — and seeds/batch dirs). You write passes; the harness +pattern-matches them into each variant's graph, checks correctness vs eager, and measures +end-to-end speedup. Your deliverable is a final report (§8). + +This skill is self-contained: every decision below is made from each variant's `model.py` +(the computation), `weight_meta.py` (shapes/dtypes), and the evaluation results. It does not +depend on any external helper script. It does, at two points, rely on *capturing the real +graph and running the real matcher/kernel* (to confirm a pattern matches and a kernel is +correct before spending a GPU eval) — that is a local, GPU-free check you can run yourself. +The `passnet-feedback` scripts package that check for convenience; use them if present, but +the flow never requires them and you must have a hand fallback when they can't capture a +graph. + +--- + +## 1. The score is the strategy + +Four metrics are reported; the **AS Score** (an ES-weighted geometric mean of per-variant +"rectified speedups") is the headline, and the other three (G-Mean Speedup, Correctness, +fast_1) move with it. Optimize the AS Score. Two facts about it drive every decision below: + +**The rectification cliff.** Per variant, your outcome is worth roughly: + +| outcome | worth (rectified speedup) | +|-------------------------------------------|---------------------------| +| no pass matched, or a crash/runtime error | **~0.1** | +| matched but numerically wrong | **~0.15** | +| matched + correct, end-to-end speedup `s` | **~s** (even when s < 1) | + +So a **matched, correct pass that runs at 0.8× scores ~8× better than not matching (0.1).** +Slow-but-correct is a *good* outcome; not-matching is a disaster. + +**The geometric-mean trap.** The sample score is a geometric mean across **all** variants. +One variant at 0.1 drags the whole sample down hard — a single red variant outweighs a hero +result on another. **Consistency across every variant beats a peak on one.** + +Consequences this skill enforces (internalize these — they explain every rule that follows): +1. **Never finish with an empty or non-matching `pass_dir`.** A floor pass is mandatory. + Matching even a *free-in-eager* region for a sub-1.0 score is far better than 0.1 — when + that's the only region you can bind to, take it (a constrained floor of ~0.5 still beats + 0.1 by ~5×). +2. **Never ship a change that turns a previously-green variant red.** Protect the gmean. +3. **Prefer ONE pattern that covers every variant** over per-variant passes — fewer ways to + leave a variant unmatched, and one evaluation validates them all. The way to do this is to + write patterns with NO shape/scalar literals that differ across variants. +4. **You do NOT have to cover every observable output.** A single matching pass already + modifies the graph (no 0.1); any output your passes don't produce simply runs in eager and + still scores. So target the **highest-value single region** rather than trying to absorb + the whole graph. (Tiny diffs may appear on outputs your kernel never touched — e.g. eager + dtype-cast nondeterminism — and they pass baseline tolerance; don't chase them.) + +(The exact tolerance weights and scoring formula are mechanics — see passnet-skill / its +references. You don't need them to make decisions; you need the cliff and the trap. + +When the harness reports an **authoritative aggregated score**, trust THAT over any estimate. +A variant is "correct" at the dtype's *baseline* tolerance — not the strict bit-exact column. +Read the end-to-end speedup, the success/failed status, and baseline-tolerance correctness; +do not be alarmed by a nonzero diff that is still within baseline tolerance.) + +--- + +## 2. "Can I even win?" — the cost gate + +The compiled path pays a **fixed per-call tax** (graph-guard + interpretation + wrapper + +launch overhead) that does not exist in eager. It is small on tiny graphs and grows with +graph size; each Triton kernel launch adds more on top. You don't need the exact numbers — +you need the rule they imply: + +> **You only win when the eager work you absorb into your kernel(s) per launch exceeds the +> tax you add.** Replacing one cheap op on a small graph loses; absorbing many ops, or +> replacing genuinely heavy memory-bound work, wins. + +Calibrate the real tax for *this machine and graph* from your first evaluation (compare the +compiled median against eager minus the work you removed) and reuse it when estimating later +regions. The detailed performance model lives in **passnet-triton-opt**; bring just enough +intuition here to triage floor-vs-ambitious in §3. + +--- + +## 2.2 The one heavy-op exception — non-overlapping conv reduces to a patch matmul + +The right question is never the op's *name*, it is its *regime*. A vendor kernel +(cuDNN/cuBLAS) is near-optimal only for (a) the shapes/parameters it has tuned kernels for and +(b) the *general* form of its op class. You can beat it only when the op's concrete parameters +put it OUTSIDE that efficient regime AND make it equal to a cheaper computation you can express +as a legal Triton kernel. Otherwise you are competing with a vendor library on its home turf and +will lose. Derive this from the parameters in front of you — do not memorize a per-op verdict. + +The clearest checkable win, and the one to actively look for: + +- **A convolution whose windows do not overlap** — `stride == kernel_size` along each spatial + dim, `kernel_size > 1`, `padding == 0`. Each output reads a disjoint input patch (no input is + reused), so the conv is *exactly* a dense matmul of flattened patches by the flattened weight + (`patches @ W.reshape(Co,-1).T` plus bias). cuDNN does not specialize for this — it runs its + general convolution path (implicit-GEMM / im2col / Winograd) built for the overlapping case, + which for an odd kernel/channel shape can be far off a well-shaped GEMM. A `tl.dot` + patch-gather kernel that also folds the flatten/transpose/tail into its epilogue can win by a + wide margin. Confirm `stride == kernel_size` and `kernel_size > 1` from the parameters. + +Why the *same* rewrite usually LOSES on the other conv/matmul shapes — i.e. why the reliable +default is "leave the op in aten and fuse its elementwise tail": + +- **1×1 conv** already *is* a pointwise matmul (PyTorch lowers it to a cuBLAS GEMM); there is no + general-conv overhead to remove, so a hand `tl.dot` competes head-to-head with cuBLAS. +- **overlapping conv** (`stride < kernel_size`, e.g. 3×3 s1): writing it as a matmul needs + im2col, which materializes a ~kernel_volume× larger matrix — the exact memory blow-up + cuDNN's implicit-GEMM/Winograd avoids. You pay more, not less. +- **depthwise / grouped conv**: memory-bound with specialized vendor kernels; a hand kernel + rarely wins and grouped indexing is bug-prone. +- **large/general dense matmul**: cuBLAS's sweet spot. + +None of this is an absolute rule — **the completed eval is always the arbiter.** The lines above +are strong *priors*, not prohibitions: if the concrete parameters give you a specific reason to +believe a case sits outside the vendor's efficient regime (e.g. a tiny/degenerate matmul, or a +memory-bound op you can fuse with a long tail into one kernel), you may try it — but only AFTER +banking the tail/floor pass, and you KEEP it only if a completed eval beats that floor. Do not +spend budget rewriting an op that is squarely in the vendor's sweet spot on a hunch, and do not +skip the non-overlapping case just because no rule forbids it. + +Two matchability facts (independent of the above): `conv1d/2d/3d`, `matmul`, `F.linear` are +C-bound functions the pattern tracer binds to fine — "it's a big cuDNN op" never means the +pattern can't match. + +Discipline (unchanged and mandatory): **always ship the tail/floor pass FIRST** (S3); only then +attempt a rewrite as an additional ambitious region; REVERT (S5/GATE e) if a completed eval does +not beat the banked floor. The rewrite is upside insurance layered on top of the floor, never a +replacement for it. + +--- + +## 2.1 Round triage gate — choose likely-amortizable samples first + +When planning a multi-sample round, do not dispatch workers randomly across small graphs. +Spend a cheap triage pass first and label each candidate sample: + +- `likely_amortizable`: at least one legal single-output region absorbs substantial real + compute per launch: a large-output broadcast/binary op, a sizeable reduction plus its + arithmetic, a long elementwise/norm chain over enough elements, or a compute tail that can + be fused without reimplementing a vendor-optimal heavy op. + Judge this from candidate regions, not only whole-sample labels: a sample with global + returned values or producer fanout can still be likely amortizable when a downstream + consumer-side region leaves the shared producer outside the pattern and returns one legal + output. +- `downside_cap_only`: only tiny regions, layout/view/reshape-only work, returned views, + split/fanout layout tails, or cheap consumers behind boundaries that are not worth + repairing. These can still justify a floor pass to avoid 0.1, but they should not consume + scarce worker slots in a performance-strategy round unless you need calibration examples. +- `blocked_by_contract`: every promising compute region would require multiple externally + observable outputs, hiding a shared producer with outside consumers, RNG/data-dependent + work, or a call form that cannot be recovered by exact/manual-FX matching. These are + pattern-fusion lessons, not performance targets. +- `unstable_eval_risk`: previous or local evidence suggests evaluator timing instability, + timeout risk, or many variants with heavy compile cost. Track this separately from + numeric, no-match, and unauthorized-operator failures; it may justify rerunning an + unchanged confirmed state, not rewriting the kernel. + +Dispatch priority for performance validation: +1. `likely_amortizable` long memory-bound elementwise or broadcast-affine-activation chains: + layout/broadcast producers plus multiple affine, activation, residual, or normalization + consumers that can be contained in one single-output region. Repeated round evidence shows + these are the best first allocation for performance-strategy validation because they can + save multiple memory passes per Triton launch. +2. Clean normalization-affine compute prefixes and concat-to-normalization-affine prefixes: + allocate formal workers to them after verifier/probe success, but record mixed outcomes + separately. The same static family can produce strong wins, near-eager correct states, or + a reverted floor depending on size, branch structure, dtype boundaries, and whether the + profitable prefix is truly the actual winning region. +3. Local-region candidates recovered by a backward slice from a promising output. Allow + shared, returned, vendor-heavy, or otherwise risky producers to stay outside the pattern + as inputs when the downstream region itself has one output and absorbs enough compute. + Repeated evidence shows this avoids over-pruning while preserving the single-output + contract. +4. Other `likely_amortizable` samples with all key anchors matchable. +5. `likely_amortizable` but one boundary uncertain, if a cheap pre-flight can settle it. +6. A small number of reduction-only or vendor-tail controls. Allocate these deliberately, + not as filler: repeated results show they often end correct-but-slow unless a larger + compute-containing region is both legal and numerically clean. +7. Defer `blocked_by_contract` unless the round is specifically about matcher behavior. + +For batch planning, keep separate fields for `triage_family`, `worker_confirmed_family`, +and `actual_winning_region_family`. Do not credit a successful larger replacement to the +initial triage family when the completed evaluation shows the winning region was really a +different family, such as a compute-plus-normalization prefix, a layout materialization tail, +or a vendor-heavy replacement. Use family-level completed-evaluation counts to adjust the +next batch allocation: wins above eager, correct-but-slow, numeric-preflight blocks, +correctness regressions on larger regions, and larger-region attempts. + +This gate is benchmark-general: the goal is to spend workers where absorbed eager work can +plausibly pay for wrapper/framework/launch tax, then use completed evaluations to confirm. + +--- + +## 3. The decision flow + +Follow this top to bottom. Gates are marked **(a)–(e)**; sub-skill hand-offs are marked +**→→ CALL**. + +``` +S0. SET UP (mechanics — see passnet-skill) + - Read graph_list + every variant's model.py and weight_meta.py. + - In service mode, confirm service health and verify uploaded pass files with /files before any /evaluate call. + - In Codex managed sandboxes, try bounded normal localhost curl first; if it fails, retry once and then use the approved escalated localhost path if available. Do not interpret sandboxed connection failures as service downtime. + - Clear stale files from pass_dir (it is imported as a package; leftovers shadow yours). + +S1. CHARACTERIZE (no GPU; read the graphs, then verify matchability with the real matcher) + - For each variant list: ops in their EXACT written call form, tensor shapes, dtypes. + - Mark each node matchable / unmatchable for normal callable patterns (the rules are in + passnet-pattern-fusion). RNG/data-dependent nodes are hard walls. Keyword-form Python + functional calls are walls for callable patterns, but a high-value single-output region + may be recoverable with an exact manual FX `GraphModule` pattern; require real matcher + proof before treating it as legal. In-place / augmented-assign ops (a node written as + `x op= c`, or a call flagged in-place) often normalize to a different form than your + pattern produces; bind to neighbours unless passnet-pattern-fusion proves an exact + pattern can match and the replacement preserves the eager mutation semantics across + repeated calls. + - Partition the graph into maximal matchable REGIONS, each with exactly ONE value + consumed outside it. Eyeballing model.py tells you the candidate regions, but whether a + given op actually matches is often NOT visible by eye — confirm it with the real matcher + (capture the graph + run the subgraph matcher, see S3 pre-flight) before you commit the + strategy to a region. + - Mark region economics before authoring: compute-heavy and large-output regions are + likely amortizable; layout/view/reshape-only regions, returned views, and split/fanout + layout tails are downside caps unless they bridge compute or remove a real + materialization. + - Classify the sample against the archetype table (§4) — this names your line of attack. + - Cross-variant literal check: do shape/scalar LITERALS differ across variants? If yes, + a shared pattern must avoid those literals (or you need per-shape passes). Every + variant must end with >=1 matching pass. + +S2. CEILING TRIAGE ===== GATE (a): FLOOR vs AMBITIOUS ===== + Estimate, per region, whether absorbed eager work clearly exceeds the tax (§2). + - Total eager work tiny / no region absorbs more than the tax + -> OVERHEAD-BOUND. Target = a safe FLOOR only. (Ceiling < 1; do not chase >1.) + - The only legal matchable regions are layout/view/reshape/split-fanout floors + -> DOWNSIDE CAP. Evaluate the safest correct state if needed, but do not tune it + blindly; stop after a completed correct eval unless a larger compute-containing + region becomes legal. + - The dominant cost is a heavy conv/matmul in the vendor sweet spot (1×1, depthwise/grouped, + overlapping conv, or a general dense matmul) with no worthwhile cheap tail + -> DEFAULT to not reimplementing it: floor on an adjacent node; accept ~1.0. + Precondition: you checked §2.2. The clearest upside case (below) is a non-overlapping conv + (`stride == kernel_size`, `kernel_size > 1`, `padding == 0`); the others are strong-prior + walls, not absolute — a completed eval may still justify a rewrite in a specific + off-regime shape. + - There is a fusible region whose absorbed work clearly beats the tax, OR a heavy op that + §2.2 shows sits outside the vendor's regime and equals a cheaper computation + -> REAL UPSIDE. Ambitious path enabled. + +S3. SHIP THE FLOOR FIRST (cap the downside before taking any risk) + Pick the single most trivially-correct matchable region — one compute node you cannot get + numerically wrong. This caps EVERY variant at "matched + correct." + →→ CALL passnet-pattern-fusion — author the floor pass (pattern + replacement + a single + kernel; use the shared-dispatch layout from the start even for one pass). + - PRE-FLIGHT (no GPU) — MANDATORY before every eval (this is the single biggest lever for + not wasting evals). Check two things: + (1) MATCH: re-read your pattern against each variant's model.py — exactly one output? + exact call form (method vs function, positional vs kwargs, no in-place op inside)? + literals equal across the variants it must cover? every inside-node consumed only + inside? Confirm it would match EVERY variant. + (2) NUMERICS: run your kernel once on inputs of the right shape/dtype and compare to + the eager computation at the dtype's baseline tolerance. Index-arithmetic kernels + (copies, permutes, strided gathers) are the #1 source of silent wrong answers — + always numerically check them before an eval. + The reliable way to do BOTH checks without a GPU eval is to reproduce the harness + locally: capture each variant's real graph (compile the model with a graph-capturing + backend, or trace it) and run the actual subgraph matcher with your pattern against it, + then run your kernel on correctly-shaped/dtyped inputs and compare to eager. The + passnet-feedback checker packages exactly this — use it if available. But it is only a + convenience: it can itself fail to capture some graphs (e.g. ones using advanced integer- + tensor indexing or other data-dependent ops). When it fails, fall back to capturing the + graph yourself the same way and running the matcher/numeric compare directly — do NOT + skip the pre-flight and "just try an eval." Reading model.py by eye is the weakest form + of this check and misses non-obvious unmatchable forms; prefer the real matcher. + - EVALUATE (GPU eval #1). ===== GATE (e) applies after every eval ===== + * service upload unconfirmed -> do NOT evaluate; fix upload/listing first. + * empty/interrupted/non-JSON eval response -> no metrics exist; retry or fix service access/state before making strategy conclusions. + * not matched / wrong -> fix MATCH or NUMERICS and re-eval; a floor that doesn't + match is worse than useless. + * matched + correct -> snapshot this pass_dir as BEST-KNOWN-GOOD; calibrate the tax + from this eval's numbers. + - If GATE (a) said FLOOR-only -> go S7. Else -> S4. + +S4. BUILD UPSIDE — one region at a time + Choose the single highest-expected-gain region not yet addressed. + →→ CALL passnet-pattern-fusion — pick which ops to absorb into the region, author the + pattern, and (for >=2 passes) wire the shared-dispatch architecture. + GATE (b): any time you author or repair a pattern, or decide what to fuse, you are + in passnet-pattern-fusion — not here, not passnet-triton-opt. + - PRE-FLIGHT (no GPU), MANDATORY — same two checks as S3 (match every variant + numeric + sanity at baseline tolerance). Fix any match/numeric problem BEFORE spending a GPU eval. + - EVALUATE (GPU eval #k). Read per-variant status + speedups + the authoritative score. + +S5. ===== GATE (e): REVERT or KEEP ===== (run after EVERY eval) + - Authoritative score < BEST-KNOWN-GOOD, OR any variant that was green is now red + -> REVERT to BEST-KNOWN-GOOD immediately. Do not stack a new change on a worse + base. Diagnose from the log, then try a DIFFERENT single change or abandon the + region. + - Otherwise -> this becomes the new BEST-KNOWN-GOOD. Continue. + +S6. ===== GATE (c): matched + correct but SLOW? ===== + - A region matches and is correct but its speedup is below the ceiling, OR correctness + misses by a SMALL numeric margin (a dtype-tolerance miss, not a gross bug) + →→ CALL passnet-triton-opt — perf tuning and/or numeric-fidelity fix. Change ONE + class at a time (perf XOR numeric). Re-eval -> back to S5. + - A completed eval is matched and correct but slower than eager, and the region is already + maximal or only layout/view/fanout work remains + -> OVERHEAD-BOUND STOP. Keep the best completed state as the floor/downside cap; + do not spend evaluations on block-size sweeps or tiny independent kernels. + - A wider legal attempt completed but scored lower than the floor or narrow state + -> REVERT and STOP that line unless a different likely-amortizable region remains. + - A broad manual-FX attempt is matched and correct but unstable or slower after absorbing + an in-place node + -> REVERT to the best state and consider a narrower downstream region that leaves + the side effect in eager. + - Matched, correct, at/near ceiling -> this region is done. + + ===== GATE (d): STOP? ===== Stop when ANY of: + * best score >= ~0.9 x your estimated ceiling, OR + * GPU-eval budget exhausted AND a matched+correct pass is in place, OR + * the last 2 evals produced no improvement, OR + * no un-addressed region has expected gain greater than the tax. + Otherwise pick the next region -> S4. + CEILING SANITY: if a dominant conv is the reason your ceiling is ~1.0, first confirm (§2.2) + it is NOT a non-overlapping conv (`stride == kernel_size`, `kernel_size > 1`, `padding == 0`) + — that one case has a much higher ceiling via a patch-matmul rewrite. Do NOT "preserve + budget" by stopping early when that specific reformulation remains unattempted and evals + remain. (For 1×1 / depthwise / general matmul a ~1.0 ceiling IS correct — don't chase them.) + +S7. FINAL REPORT (§8). Make sure BEST-KNOWN-GOOD is what is on disk. +``` + +Compact tree: + +``` +read sample -> characterize -> ceiling triage (a) + overhead-bound / vendor-heavy -> floor (b) -> eval -> report + real upside -> floor (b) -> eval -> [loop: pick region (b) -> eval + -> revert-or-keep (e) + -> if slow/near-miss tune (c) + -> stop? (d)] -> report +``` + +--- + +## 4. Graph-shape archetypes → line of attack (operator-AGNOSTIC) + +Classify by **dataflow structure, tensor-size regime, and matchability — never by which +operator it is**. (The right kernel is derived from the graph in front of you; this table +only tells you *where the upside lives* and *what to try first*.) + +| structural shape | size regime | line of attack | rough ceiling | +|---|---|---|---| +| a chain of cheap elementwise/normalize ops feeding ONE output | any | fuse the maximal matchable region into one kernel | good if tensors large; ~1.0 if tiny | +| a reduction over a large axis, then arithmetic on the result | ms-scale tensors | one fused reduction kernel producing the final output | among the highest (often >1.5×) | +| a broadcast/expand feeding a binary op that WRITES a large output | large output | one shape-aware kernel; absorb the layout ops | strong even alone | +| a heavy conv/matmul + a cheap tail — 1×1 conv, depthwise/grouped conv, overlapping conv, general dense matmul | any | default: leave the heavy op in the graph; fuse only the tail. Rewriting the op is usually a loss (vendor sweet spot) — try only if parameters suggest off-regime AND a completed eval beats the floor | bounded by tail size | +| a non-overlapping conv (`stride == kernel_size`, `kernel_size > 1`, `padding == 0`) | any | reformulate as a `tl.dot` patch-gather kernel + fused tail; the conv is the region ANCHOR, not a wall. Ship a tail floor first | can be large | +| tiny total eager time (a couple of cheap ops) | small | overhead-bound: ship a floor, accept < 1 | < 1 | +| mostly layout ops (view/reshape/permute/transpose) + 1 compute op | any | replace the compute op only; absorb adjacent layout for free | modest | +| the only matchable region is a layout op whose result is RETURNED directly (a view eager never materializes) | any | replacing it FORCES a copy eager skips, so it scores < 1 — but it still matches; take it ONLY as a last-resort floor to escape 0.1 | < 1, but >> 0.1 | +| split/fanout layout region with multiple externally observable consumers | any | single-output floor or separate consumer-side passes only; do not hide the shared producer | < 1 unless each pass absorbs real compute | +| graph contains RNG / data-dependent / unrecovered call-form / in-place nodes | any | those nodes are walls by default; try exact manual FX only for valuable single-output call-form blockers, otherwise fuse around them | bounded by what's left | + +Rules for using this table: +- **Matchability overrides structure.** A graph can look exactly like a high-upside row (a + reduction-then-arithmetic chain, say) yet score < 1, because the op that would anchor the + fusion is in a form a normal callable pattern cannot reproduce, or because an in-place, + RNG, or data-dependent node severs the region. So BEFORE you commit to an ambitious row, + confirm the anchor op actually matches. For kwargs-form Python functional blockers, first + try the normal callable form; if it near-misses and the region is valuable, use + passnet-pattern-fusion's manual FX pattern path and verify it with the real matcher. + **Verify matchability with the real matcher** (the passnet-feedback checker, or a graph + capture you do yourself per S1) before building on a region — not just before the eval. +- Pick the row that matches the **structure** AND whose anchor op is matchable, then hand the + region to **passnet-pattern-fusion** to turn into a concrete pattern + kernel. +- Layout-only ops are ~free in eager: absorbing them pays only when it lets you bridge two + compute ops into one kernel or erase a materialization — not as a goal in itself. The one + exception is the last-resort floor above: when EVERY compute op is unmatchable, materializing + a returned view is the only way to modify the graph, and a ~0.5 beats a 0.1. +- If two rows apply, prefer the one with the larger absorbed-work-per-launch. +- Your kernel can be faster than eager in isolation yet the sample still lands near 1.0, + because an unmatchable neighbour you couldn't absorb stays in BOTH paths. The ceiling is set + by absorbed work, not total work — estimate it from what you can actually bind to. + +--- + +## 5. Dead ends (DON'T) and disciplines (DO) + +**DON'T** (each line is a path that scores badly — avoid it): +- **Finish with an empty or non-matching `pass_dir`.** That's the 0.1 floor; a slow correct + pass is ~8× better. +- **Return more than one value from a pattern.** Multi-output crashes the variant -> 0.1. + If two intermediates are observable, that's two separate passes (passnet-pattern-fusion). +- **Add a second pass without the shared-dispatch architecture.** With >=2 passes the harness + silently drops any pass whose replacement function differs -> variants stop matching. +- **Re-implement a vendor-optimal heavy op in a custom kernel that does the SAME work** — a + general dense matmul, a 1×1 conv (already a pointwise GEMM), a depthwise/grouped conv, or an + overlapping conv, when it sits in the vendor's efficient regime. You will lose to the library; + fuse its cheap tail instead. (The clearest exception is not "the same work": a non-overlapping + conv — `stride == kernel_size`, `kernel_size > 1`, `padding == 0` — is algebraically a dense + patch matmul while cuDNN runs a heavier general path, so a `tl.dot` patch-gather rewrite is a + real algorithm change worth pursuing; see §2.2. This red line is a strong prior — a completed + eval can still justify a rewrite for a provably off-regime shape. Also: never conclude a heavy + op is unmatchable — it is a C-bound function that binds fine.) +- **Chase >1.0 on an overhead-bound sample.** The tax caps the ceiling below 1 — take the + floor and move on. +- **Tune a correct-but-slow tiny/layout-only/fanout floor by reflex.** First ask whether a + larger legal compute-containing region exists; if not, stop with the floor. +- **Make multi-change rewrites between evals.** Regressions become un-diagnosable. One + change-class per eval. +- **Build on top of a change that lowered the score.** Revert to BEST-KNOWN-GOOD first. +- **Put RNG / data-dependent / unrecovered call-form nodes inside a pattern.** They are + region boundaries; including them without exact matcher proof means no match. +- **Evade the anti-cheat** (operator laundering, dispatch tricks, returning precomputed eager + results). It's classified as hacking, scores 0 on the leaderboard, and is usually slower + than eager anyway. +- **Spend a GPU eval to discover a match failure** you could have caught by running the real + matcher against the captured graph, or a numeric bug you could have caught by running the + kernel once on correctly-shaped inputs. +- **Mis-judge correctness from a strict bit-exact comparison.** A variant counts as correct + at its dtype's *baseline* tolerance, which is loose for fp16/bf16. A small nonzero diff — + especially on an output your kernel never produced — is usually fine. Read the authoritative + speedup/score, not the strict-equal column. +- **Abandon a region as "infeasible" because your first kernel was wrong.** A wrong result is + a numeric bug to fix (one change-class), not proof the strategy can't work. + +**DO** (process discipline that the trajectories show pays off): +- **Ship a trivially-correct floor pass first**, then build upside on a known-good base. +- **One change-class per iteration** (match | numeric | perf) so every eval is attributable. +- **Track a BEST-KNOWN-GOOD snapshot** and revert to it on any regression. In service mode + keep a local copy of every uploaded version. +- **Write patterns with no shape literals when variants differ only by shape** — one pass + then covers all variants and one eval validates them. +- **Calibrate the per-call tax from your first eval** and reuse it in ceiling estimates. +- **Maximize ops absorbed per launch** rather than shaving an already-tiny kernel. +- **Recognize overhead-bound samples early and accept a sub-1.0 floor.** +- **Classify evaluator timing instability separately** from no-match, numeric, + unauthorized-operator, timeout, and OOM failures. If the pass pre-flights cleanly and + successful variants are correct, an unchanged rerun can be a valid use of budget. +- **Abandon a fundamentally-losing line of attack early** (don't sink evals into fighting a + vendor kernel or a region whose absorbed work can't beat the tax) — but distinguish "wrong + strategy" from "right strategy, buggy kernel": fix the bug before abandoning. +- **Read the authoritative result.** Use end-to-end speedup + success status + baseline- + tolerance correctness as ground truth; treat any other estimate as a hint, not a verdict. + +--- + +## 6. GPU-eval budget & iteration discipline + +GPU evaluations cost minutes (many trials × many variants) and there is a hard per-eval +timeout. Treat them as the scarce resource and plan around them: + +- **Budget ~6–8 GPU evals per sample.** If you're past that without progress, ship + BEST-KNOWN-GOOD and stop. +- **Eval #1 is always the floor pass** — it buys downside insurance and calibrates the tax. + Never spend eval #1 on an ambitious multi-pass attempt. +- **Pre-flight before every eval is mandatory** (both checks in S3: match-every-variant + + numeric sanity at baseline tolerance — by hand or with the optional checker). A GPU eval + that only learns "didn't match" or "kernel was wrong" is wasted; the budget assumes you + never spend one that way. Done well, this drives wasted evals toward zero. +- **One change-class per eval** — so each eval answers exactly one question. +- **Many-variant samples (tens to 100+ graphs) risk timeout.** There, minimize compile cost: + no autotune, a single fixed kernel config, the fewest passes, and a small fixed set of + compile-time specializations. (The *how* is passnet-triton-opt; the *decision to do it* is + yours, here, based on variant count.) +- **Snapshot every evaluated pass_dir state** so reverting is instant. + +--- + +## 7. Generalization & anti-copy + +The sample you are given will differ from anything you've seen. **Derive the optimization +from the dataflow in `model.py` and the shapes/dtypes in `weight_meta.py` in front of you** — +not from a remembered "for operator X do Y" recipe. If you catch yourself recalling a +per-operator formula, stop and re-derive it from the actual graph; the held-out set is built +to punish memorized answers. This skill deliberately gives you *decision structure* (where to +start, what to try, what to avoid), and leaves the operator-specific kernel math to be worked +out per problem in passnet-pattern-fusion / passnet-triton-opt. + +--- + +## 8. Final report + +End your work with: +- **Final AS Score** (from the last evaluation) and whether the pass matched. +- **Per-variant summary**: dtype → matched? correct? speedup. +- **Pass files** you created and a one-paragraph strategy description. +- If below the ceiling you estimated: the specific blocker (overhead-bound, vendor op, + numeric limit, timeout, …). +- Do NOT write result files to disk; your reply text is the result. + +--- + +## 9. Where things live (boundary map) + +| skill | owns | +|---|---| +| **passnet-orchestrate** (this) | the decisions: where to start, what to try, dead ends, floor-vs-ambitious, revert/stop, eval budget | +| **passnet-pattern-fusion** | authoring patterns that match, choosing fusion regions, multi-pass shared-dispatch wiring, kernel templates | +| **passnet-triton-opt** | making a matched kernel fast (block/grid/warps/autotune/launch) and numerically faithful per dtype | +| **passnet-skill** | mechanics: fetching the problem, the pass-file format, how evaluation is invoked | +| **passnet-feedback** (optional) | scripts that give the same analysis/pre-flight/log-parsing faster — never required by this flow | + +This skill is the **strategy** half and **passnet-skill** is the **mechanics** half; use them +together (this one to decide, that one for how to fetch/format/evaluate). If you only need to +know *how* to do a mechanical step, go straight to passnet-skill; for *what to do and in what +order*, stay here. diff --git a/skills/task-oriented/Passnet/skills/passnet-pattern-fusion/SKILL.md b/skills/task-oriented/Passnet/skills/passnet-pattern-fusion/SKILL.md new file mode 100644 index 000000000..4b0fb9eef --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-pattern-fusion/SKILL.md @@ -0,0 +1,333 @@ +--- +name: passnet-pattern-fusion +description: > + Author PassNet pass files: write patterns that actually MATCH the FX graph, pick fusion + regions (which ops to absorb into one Triton kernel), and structure multi-pass + submissions with the mandatory shared-dispatch architecture. Use when creating or fixing + pattern()/replacement_args()/replacement_func() files, when a pass "failed to match", + or when deciding how to fuse multiple kernels/ops. +--- + +A pass file teaches the compiler: "this subgraph → my Triton kernel". Getting the PATTERN +right is binary (match or 0.1); getting the FUSION REGION right decides the speedup. +Kernel implementation details live in `references/kernel-templates.md`; numeric-fidelity +recipes in passnet-triton-opt. + +## 1. Pass file contract + +```python +import torch # 'import triton' etc. as needed +def pattern(a, b, ...): # subgraph to match — torch IR, mirrors model.py + ... + return out # EXACTLY ONE returned value (tuple-of-1 ok) +def replacement_args(a, b, ...): # pure arg shuffle, FX-traced — no math/.shape here! + return (a, b, "route_name") # constants (route strings/scalars) may be appended +def replacement_func(): # returns the SAME module-level function every call + return dispatch_wrapper +``` +Plus `pass_dir/sorted_output_pass_rule_names.json`: `["PassA", "PassB"]` — file stems, +priority order, every pass you want loaded MUST be listed. + +## 2. The 10 hard rules (violating any ⇒ no match / crash / dropped pass) + +1. **ONE output.** A pattern returning ≥2 values crashes the variant (harness asserts). + If two intermediates are observable outside the region, write two passes that each end + at one of them — never one pattern returning both. +2. **Mirror `model.py`'s exact call forms.** The target graph keeps the source form; your + pattern is traced with normalization (see matchability table below). Same op spelled + differently = no match (`torch.relu` ≠ `torch.nn.functional.relu`; `x.transpose(...)` + (method) ≠ `torch.transpose(x, ...)`). +3. **Literals must be exactly equal**: eps `1e-05` ≠ `1e-04`, dims, shape lists, scale + constants, `device(type='cuda')` objects in factory calls. Shape literals differing + across variants ⇒ that node can't be in a shared pattern (write per-shape passes or + exclude the node). +4. **Containment**: every node inside the pattern except the returned one must be consumed + ONLY inside the pattern. If a producer feeds both your target consumer and any outside + or returned value, do not hide that producer inside a larger pattern. End the region at + the consumer/tail output or write separate one-output passes. If `model.py` returns an + intermediate or feeds it to a node outside your region, end the region there (it becomes + the pattern's single output) or exclude it. When an external producer has multiple + observable consumers, prefer a one-output consumer-side floor pass or split regions + over pulling that producer into a fused region. + This is the single-output contract in practice: a producer whose value is externally + visible or shared across a fanout cannot disappear inside your pass unless the one value + returned by the pattern preserves every outside observation. Do not "simplify" by + recomputing only the consumer you care about while hiding the shared producer; that makes + the graph semantically incomplete or non-containment-safe. +5. **No `tmp = None`** statements, no prints, no asserts inside `pattern`. Pure alias + lines in model.py (`tmp_0 = in_0`) create NO graph node either — skip them and use the + original name's value flow. +6. **`replacement_args` is FX-traced**: only reorder/drop arguments and append literal + constants. No arithmetic, no `.shape`, no conditionals on tensors. +7. **`replacement_func()` must be stable** (module-level function, `f() is f()`), and with + 2+ pass files ALL must return the SAME object via the shared-dispatch module (§4). +8. **No RNG ops** (`torch.rand*`, dropout `training=True`) inside patterns; eval-mode + dropout (`training=False`) is identity and safe to absorb. +9. **AST restrictions** (file-wide): never import `torch.nn(/functional)`, `torch.ops`, + `torch.autograd`, never alias them; in pattern bodies use full dotted + `torch.nn.functional.xxx(...)`. Outside `pattern`/`replacement_args`, only the + whitelisted `torch.empty/zeros/ones/full[_like]/as_tensor` calls. +10. **Wrapper-runtime restrictions** (poison dispatch during warmup): inside the kernel + wrapper, only allocator ops + metadata (`.shape/.stride()/.numel()/.dim()/.device/ + .dtype/.data_ptr()`) + `.to()` + Triton launches. NO `.contiguous()`, no tensor math. + +## 3. Matchability — what a pattern node can bind to + +For a normal callable `def pattern(...)`, the pattern is traced with `ForceArgsTracer`; the +target (dynamo) graph is not. Per node: + +| node kind in model.py | pattern is normalized? | matchable when... | pattern should write... | +|---|---|---|---| +| method call `x.view(...)`, `x.mean(dim=-2, keepdim=True)`, `x.to(dtype=...)`, `x.softmax(dim=-1)` | NO | always | mirror EXACTLY (incl. kwargs spelling) | +| C-bound function: `torch.conv2d`, `torch.matmul`, `torch.cat`, `torch.sigmoid`, `torch.arange`, `torch.sum`, `F.linear`, `F.gelu`, `operator.+ - * /` | NO (no inspect signature) | always | mirror EXACTLY (incl. kwargs spelling) | +| Python-def `F.*`: `relu`, `silu`, `softmax`, `dropout`, `layer_norm`, `batch_norm`, `embedding`, `interpolate`, `adaptive_avg_pool2d`, `pad`, `normalize`, `max_pool2d`... | YES → full positional, defaults filled | ONLY if model.py call is full-positional with ALL params spelled out | the same call (any style — it normalizes) | + +Examples from real graphs: +- `torch.nn.functional.batch_norm(x, m, v, w, b, False, 0.1, 1e-05)` → matchable ✓ +- `torch.nn.functional.dropout(x, 0.1, False, False)` → matchable ✓ (and it's identity!) +- `torch.nn.functional.relu(x, inplace = False)` → **kwargs on a Python-def op: NOT matchable** ✗ +- `torch.nn.functional.softmax(x, 2, _stacklevel = 5)` → ✗ (and `_stacklevel` precedes + `dtype` in the signature — normalization can't reproduce it) +- `x.softmax(dim=-1)`, `x.mean(dim = -2, keepdim = True)` → method form: matchable ✓ +- `torch.cat([a, b], dim = 2)` → C-bound with kwargs: matchable (mirror the kwargs) ✓ +- `torch.nn.functional.gelu(x)` / `F.gelu(x, approximate='none')` → C-bound: matchable ✓ + +Treat kwargs-form Python-level activations and functional calls as suspect boundaries +until `check_pattern.py` proves they match. For small or low-value regions, if they block +matching, start the region after that node or end just before it instead of spending +evaluations on near-miss patterns. + +**Manual FX escape hatch for valuable kwargs regions.** The harness accepts `pattern` as a +`torch.fx.GraphModule`. In that case `_replace_pattern` uses `pattern.graph` directly and +does NOT run `ForceArgsTracer`, so exact kwargs-form Python functional nodes can match. +Use this only when a high-value single-output region is blocked by callable-pattern +normalization and a real pre-flight matcher confirms the manual graph matches. + +Build the graph inside a function literally named `pattern` (that function body is +AST-exempt), then replace `pattern` with the returned `GraphModule` and set +`__signature__` to the placeholders in order: + +```python +import inspect +import operator +import torch + +def pattern(): + graph = torch.fx.Graph() + x = graph.placeholder("x") + y = graph.placeholder("y") + act = graph.call_function( + torch.nn.functional.relu, args=(x,), kwargs={"inplace": True} + ) + out = graph.call_function(operator.add, args=(act, y), kwargs={}) + graph.output(out) + gm = torch.fx.GraphModule({}, graph, "ExactKwargPattern") + gm.__signature__ = inspect.Signature([ + inspect.Parameter("x", inspect.Parameter.POSITIONAL_OR_KEYWORD), + inspect.Parameter("y", inspect.Parameter.POSITIONAL_OR_KEYWORD), + ]) + return gm + +pattern = pattern() +``` + +The manual graph must still obey every other rule: one output, containment, exact targets +and literals, matching `replacement_args` parameters, and Triton-only runtime computation. +Do not use it as a speculative default; it is a repair for proven callable-pattern +normalization misses. + +Manual FX does not make side effects free. Treat `inplace=True` functional calls as +side-effectful even when the graph matches: only absorb them if the replacement mutates the +input exactly as eager would under repeated benchmark calls. If that is awkward or uncertain, +leave the in-place node in eager and fuse the downstream side-effect-free region. + +Do NOT trust memory on which ops are C-bound — VERIFY with the pre-flight checker +(`passnet-feedback/scripts/check_pattern.py`) before every GPU eval; it traces your +pattern, runs the real `SubgraphMatcher` against the real dynamo graph of every variant, +and prints near-miss diffs when a pattern doesn't match. + +## 4. Multi-pass architecture (mandatory for ≥2 pass files) + +`output_pass_replacement_func_limit: 1` keeps only ONE distinct replacement function — +passes returning a different object are SILENTLY dropped (watch `Loaded N passes` in the +log). Structure: + +`pass_dir/_shared_kernels.py` (helper module, NOT listed in the json): +```python +import torch +import triton +import triton.language as tl + +@triton.jit +def _fused_a_kernel(...): ... +@triton.jit +def _fused_b_kernel(...): ... + +def _run_a(x, w): + out = torch.empty_like(x) + ...launch _fused_a_kernel... + return out + +def _run_b(x): + ... + return out + +@torch.fx.wrap +def dispatch_wrapper(*args): + route = args[-1] + if route == "fuse_a": + return _run_a(args[0], args[1]) + if route == "fuse_b": + return _run_b(args[0]) + raise ValueError(f"unknown route {route}") +``` + +each `pass_dir/FuseA.py`: +```python +import torch +from pass_dir._shared_kernels import dispatch_wrapper + +def pattern(x, w): + ... + return out + +def replacement_args(x, w): + return (x, w, "fuse_a") + +def replacement_func(): + return dispatch_wrapper +``` + +- Works because the sample root is on `sys.path` while passes load; the shared module is + imported once → one function object → nothing is dropped. +- Keep `_shared_kernels.py` AST-clean too (Triton + whitelisted allocators only). +- A single-pass submission may simply define its wrapper in the pass file — but starting + with the shared layout costs nothing and survives growth. +- Order in the json = application order. Earlier passes consume nodes; a later pattern + overlapping an earlier match will no longer find its nodes. Order independent patterns + freely; order overlapping ones from largest to smallest region. + +## 5. Choosing fusion regions + +1. From the analyzer output, mark every node matchable/unmatchable for a normal callable + pattern, and separately mark high-value kwargs-form Python functional nodes that may be + recoverable with a manual FX `GraphModule` pattern. +2. Greedily grow regions over callable-matchable nodes along dataflow; STOP at: an + unrecovered unmatchable node, RNG, a node whose output is consumed outside the region + (unless you end the region exactly there), a conv/big-matmul you've decided to keep in aten. + **Keeping a heavy conv/matmul in aten and fusing its tail is the DEFAULT and usually right.** + The clearest case where the conv itself becomes the region ANCHOR instead of a wall: a + **non-overlapping conv** — `stride == kernel_size`, `kernel_size > 1`, `padding == 0` — which + is exactly a dense patch matmul (§ kernel-templates §13). Anchoring on a 1×1 conv (already a + cuBLAS-optimal pointwise GEMM), a depthwise/grouped conv, or a general matmul usually loses — + default to fusing their tails and leaving the op in aten (let a completed eval overrule this + if parameters look off-regime). Either way the op is C-bound and matches when you mirror its + exact positional call form; never treat "it's a big cuDNN op" as "unmatchable." + If the only blocker is callable-pattern normalization of kwargs-form Python functional + calls and the wider region has real upside, try a manual FX pattern and require + `check_pattern.py` proof before evaluation. + For fanout, draw the boundary at the first value that has multiple external consumers. + If each consumer is independently valuable, write separate one-output passes using the + shared-dispatch architecture; if the consumers are only cheap layout/view work, treat + them as floor/downside-cap candidates rather than performance targets. + When a shared or returned producer feeds a substantial downstream consumer, consider a + consumer-side region that leaves that producer outside the pattern and treats it as an + input. Do not reject the whole sample for global fanout if a local backward slice from the + chosen output still has exactly one output and absorbs enough compute to amortize a + launch. + Prefer a single-output region that absorbs several elementwise, affine, activation, + residual, or normalization consumers fed by the same broadcast/layout producer when + containment is legal. Repeated round evidence shows this family is more likely to pay the + fixed launch/wrapper cost than isolated reductions or tiny tails. + For concat-to-normalization-affine prefixes, containment can still be legal even when a + downstream activation is unmatchable or in-place: return the normalized or affine prefix + tensor as the one crossing value and leave the boundary consumer in eager. Do not pull an + in-place activation into the pattern just to make the region look complete. + When repeated duplicate compute exists, anchor the replacement on values with real users. + Userless/dead duplicate nodes can structurally match but fail replacement because the + matched return has no observable consumer. If the profitable opportunity is dead-work + elimination, treat it as a graph-rewrite edge case and prove it with the real matcher and + smoke test before spending an evaluation. +3. Each region ⇒ one pass ⇒ ideally ONE Triton kernel launch. Two kernels for one region + only when a reduction's result feeds elementwise work over a DIFFERENT axis size + (e.g. softmax over rows then matmul) — usually better: split into two passes. +4. Region inputs = tensors crossing into it (pattern parameters); region output = the one + tensor crossing out. +5. Sanity: replacing N aten ops saves roughly Σ(their eager µs); your kernel costs + ~19 µs launch + memory traffic. If the sum is < ~50 µs the region only makes sense as + the floor pass / part of something bigger. +6. Layout ops at region EDGES are usually better left outside (they're ~free in eager). + Layout ops INSIDE a region (between two compute ops) are absorbed for free via index + arithmetic. Special case worth taking even "alone": `view/unsqueeze/expand → binary op` + chains that WRITE a large output — aten's broadcast-elementwise kernels are often + 1.5–2.5× off the bandwidth roofline, so a shape-specialized Triton kernel wins well + beyond the erased python overhead (measured 2.3× on a 67M-element broadcast-sub). + Don't guess: `analyze_graph.py --bench` gives each node's real eager µs; estimate your + kernel as `bytes_moved / 1.4 TB/s` and compare. + Do not force eager metadata-only layout/view/reshape work to materialize just to make a + larger pattern. If the pattern output is a returned or externally observed view that eager + can keep cheap, materializing it in a wrapper may erase the benefit of absorbing nearby + elementwise work; treat that as a downside cap unless a completed eval proves otherwise. + +## 6. Kernel wrapper conventions (poison-proof) + +```python +def _run_fused(x, w, b): + B, N = x.shape # metadata: allowed + out = torch.empty((B, N), dtype=x.dtype, device=x.device) + grid = (triton.cdiv(B * N, 1024),) + _kernel[grid](x, w, b, out, B * N, N, + x.stride(0), x.stride(1), # pass strides — NEVER .contiguous() + BLOCK=1024) + return out +``` +- Output dtype MUST equal what eager produces (usually the input dtype; watch explicit + `.to(torch.float32)` nodes in the region). +- If the pattern output is consumed by `.view(...)` outside the region, your output's + contiguity must allow that view — return a fresh contiguous tensor of the eager + output's shape (the normal case). +- Handle every variant's shapes/dtypes with ONE wrapper: derive sizes from `.shape`, + never hardcode, unless the pattern itself is shape-literal-specific. +- Scalars you need (eps, slopes, dims) that appear as pattern literals: hardcode them in + the kernel for that pass (they're fixed by the match) or pass via `replacement_args` + constants. + +See `references/kernel-templates.md` for ready-to-adapt Triton templates: elementwise +chain (stride-aware), row reduction (softmax/layer_norm/mean/L2-norm), BN-inference +affine, bias+activation epilogue, cat+elementwise, embedding gather, matmul+epilogue +(`tl.dot`), 2-pass split reductions. + +## 7. Debug: "Pass X failed to match" + +1. Run `check_pattern.py` — it prints the dynamo node list vs your traced pattern nodes. +2. Compare node-by-node: op kind (`call_function` vs `call_method`), target identity, + arg/kwarg shapes of the WRITTEN forms, literals. +3. Most frequent causes, in order: kwargs-form Python-def F.* op included in a normal + callable pattern (drop the node for a cheap region, or switch to an exact manual FX + pattern for a valuable one); literal mismatch across variants; intermediate consumed + outside (containment); userless/dead nodes selected as the match output; wrong + method-vs-function form; `_stacklevel`-style hidden kwargs; pattern includes a + `tmp = None`-induced phantom; two returned values. If a producer has outside consumers, + exclude that producer or split at the consumer output before retrying. +4. `[PassMgrBackend] Diagnostic ... best-attempt` lines in eval logs name the first + mismatching node pair — read them. + +## 8. Round-planning containment checklist + +Before handing a region to kernel tuning, answer these with the real graph in mind: + +- Does the proposed pattern return exactly one value? +- Are all nodes hidden inside the pattern consumed only inside the pattern, except that + single returned value? +- Does any producer inside the region feed another returned value, an outside consumer, or + a later observable fanout branch? If yes, exclude it or split into one-output passes. +- Is the region mainly layout/view/reshape/split fanout work? If yes, mark it as a + downside-cap/floor candidate unless it bridges compute or avoids a real materialization. +- Does the proposed region force a returned view or metadata-only layout value to become an + allocated output? If yes, prefer a consumer-side elementwise/affine region or a narrow + floor, and require completed-evaluation proof before keeping the wider materializing pass. +- If multiple passes are needed, do they all return the same shared dispatch wrapper? + +Tentative guidance from repeated round behavior: when a larger-looking region is blocked +only by the single-output/external-fanout contract, it is usually better to stop with the +best correct floor than to add tiny independent kernels that increase launch overhead. diff --git a/skills/task-oriented/Passnet/skills/passnet-pattern-fusion/references/kernel-templates.md b/skills/task-oriented/Passnet/skills/passnet-pattern-fusion/references/kernel-templates.md new file mode 100644 index 000000000..f85e3c87f --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-pattern-fusion/references/kernel-templates.md @@ -0,0 +1,325 @@ +# Triton kernel templates for PassNet passes + +Adapt, don't copy blindly: derive sizes from tensor metadata, keep dtype-generic +(`compute fp32, store out dtype`), pass strides for any input that may be non-contiguous +(anything produced by `permute/transpose/expand/slice` upstream — check the graph). +All templates are poison-dispatch-safe (allocators + metadata only). + +## 0. Conventions + +```python +import torch +import triton +import triton.language as tl +``` +- Load pattern: `x = tl.load(ptr + offs, mask=mask, other=0.0).to(tl.float32)` +- Store pattern: `tl.store(out_ptr + offs, y.to(out_ptr.dtype.element_ty), mask=mask)` +- Grid: `grid = (triton.cdiv(n, BLOCK),)`; BLOCK 1024 default for elementwise; + `num_warps=4` default, 8 for BLOCK ≥ 2048. +- For 2D row kernels: one program per row (or per row-block), columns swept by + `tl.arange(0, BLOCK_N)` with `BLOCK_N = triton.next_power_of_2(N)` when N ≤ 16384. + +## 1. Fused elementwise chain (flat, contiguous) + +For regions like `bn-eval → add → sigmoid → mul` (any per-element math over same-shape +tensors). + +```python +@triton.jit +def _ew_chain(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr): + pid = tl.program_id(0) + offs = pid * BLOCK + tl.arange(0, BLOCK) + m = offs < n + x = tl.load(x_ptr + offs, mask=m, other=0.0).to(tl.float32) + y = tl.load(y_ptr + offs, mask=m, other=0.0).to(tl.float32) + r = x * (1.0 / (1.0 + tl.exp(-y))) # example: x * sigmoid(y) + tl.store(out_ptr + offs, r.to(out_ptr.dtype.element_ty), mask=m) + +def _run(x, y): + out = torch.empty_like(x) + n = x.numel() + _ew_chain[(triton.cdiv(n, 1024),)](x, y, out, n, BLOCK=1024) + return out +``` + +Broadcast variants: a `[C]` tensor broadcast over `[N, C]` rows → index it with +`offs % C`; a `[N, 1]` column broadcast → `offs // C`. For NCHW per-channel params +(`[C]` over `[N, C, H, W]`): channel = `(offs // (H * W)) % C` — pass `HW = H*W` and `C`. + +## 2. Stride-aware elementwise (non-contiguous input / absorbing `.contiguous()`) + +When an input comes from `permute/transpose/expand` (stride may be 0!) or you absorb a +`contiguous()` node: compute multi-dim indices from the flat output offset, then address +input by ITS strides. + +```python +@triton.jit +def _ew_strided(x_ptr, out_ptr, n, D0, D1, D2, # logical out shape + sx0, sx1, sx2, # input strides (elements) + BLOCK: tl.constexpr): + pid = tl.program_id(0) + offs = pid * BLOCK + tl.arange(0, BLOCK) + m = offs < n + i2 = offs % D2 + i1 = (offs // D2) % D1 + i0 = offs // (D1 * D2) + x = tl.load(x_ptr + i0 * sx0 + i1 * sx1 + i2 * sx2, mask=m, other=0.0).to(tl.float32) + tl.store(out_ptr + offs, x.to(out_ptr.dtype.element_ty), mask=m) +``` +Wrapper passes `x.stride(0), x.stride(1), x.stride(2)`. This is also THE template for +materializing any layout chain (`permute+contiguous`, `expand`, simple `cat` reads). + +## 3. BatchNorm-inference affine (matches `F.batch_norm(x, m, v, w, b, False, mom, eps)`) + +```python +@triton.jit +def _bn_eval(x_ptr, mean_ptr, var_ptr, w_ptr, b_ptr, out_ptr, n, C, HW, + EPS: tl.constexpr, BLOCK: tl.constexpr): + pid = tl.program_id(0) + offs = pid * BLOCK + tl.arange(0, BLOCK) + m = offs < n + c = (offs // HW) % C # NCHW; for [N, C] tensors pass HW=1 → offs % C + x = tl.load(x_ptr + offs, mask=m, other=0.0).to(tl.float32) + mu = tl.load(mean_ptr + c, mask=m, other=0.0).to(tl.float32) + var = tl.load(var_ptr + c, mask=m, other=1.0).to(tl.float32) + w = tl.load(w_ptr + c, mask=m, other=1.0).to(tl.float32) + b = tl.load(b_ptr + c, mask=m, other=0.0).to(tl.float32) + y = (x - mu) * (1.0 / tl.sqrt(var + EPS)) * w + b + tl.store(out_ptr + offs, y.to(out_ptr.dtype.element_ty), mask=m) +``` +Fuse adjacent matchable activations/residual-adds into `y` before the store. + +## 4. Row reduction — softmax (matches method-form `x.softmax(dim=-1)`) + +Eager softmax subtracts the row max — replicate EXACTLY for fp32-baseline correctness. + +```python +@triton.jit +def _softmax_rows(x_ptr, out_ptr, R, N, sx_r, BLOCK_N: tl.constexpr): + row = tl.program_id(0) + cols = tl.arange(0, BLOCK_N) + m = cols < N + x = tl.load(x_ptr + row * sx_r + cols, mask=m, other=float("-inf")).to(tl.float32) + x = x - tl.max(x, axis=0) + e = tl.exp(x) + y = e / tl.sum(e, axis=0) + tl.store(out_ptr + row * N + cols, y.to(out_ptr.dtype.element_ty), mask=m) + +def _run_softmax(x): + shp = x.shape + N = shp[-1] + R = x.numel() // N + out = torch.empty_like(x) + _softmax_rows[(R,)](x, out, R, N, x.stride(-2) if x.dim() > 1 else 0, + BLOCK_N=triton.next_power_of_2(N), num_warps=8 if N > 1024 else 4) + return out +``` +`softmax(dim=k)` for non-last dims: either pass strides and sweep that axis, or only +absorb last-dim softmax (the common case). Fuse a preceding scale (`q @ k / sqrt(d)`'s +div) or additive mask into the load expression when those nodes are in the region. + +## 5. Row reduction — LayerNorm (matches positional `F.layer_norm(x, (N,), w, b, eps)`) + +Eager computes mean and rstd in fp32: `y = (x - mean) * rsqrt(var + eps) * w + b`, var is +the BIASED variance (divide by N). + +```python +@triton.jit +def _layer_norm(x_ptr, w_ptr, b_ptr, out_ptr, N, EPS: tl.constexpr, BLOCK_N: tl.constexpr): + row = tl.program_id(0) + cols = tl.arange(0, BLOCK_N) + m = cols < N + x = tl.load(x_ptr + row * N + cols, mask=m, other=0.0).to(tl.float32) + mean = tl.sum(x, axis=0) / N + d = tl.where(m, x - mean, 0.0) + var = tl.sum(d * d, axis=0) / N + rstd = 1.0 / tl.sqrt(var + EPS) + w = tl.load(w_ptr + cols, mask=m, other=1.0).to(tl.float32) + b = tl.load(b_ptr + cols, mask=m, other=0.0).to(tl.float32) + y = (x - mean) * rstd * w + b + tl.store(out_ptr + row * N + cols, y.to(out_ptr.dtype.element_ty), mask=m) +``` +Same skeleton handles RMSNorm chains (`pow(2).mean → rsqrt → mul → mul w`), L2-normalize +(`x / x.norm(p=2, dim=-1, keepdim=True)` — method-form `.norm` matches), and +`mean(dim=-1)`/`sum(dim=-1)` if the reduced value is the region output. + +## 6. Mean over trailing spatial dims (matches `x.mean((2, 3), keepdim=True)` / +`F.adaptive_avg_pool2d(x, 1)`) + +```python +@triton.jit +def _mean_hw(x_ptr, out_ptr, HW, BLOCK: tl.constexpr): + nc = tl.program_id(0) # one program per (n, c) + offs = tl.arange(0, BLOCK) + acc = tl.zeros((BLOCK,), dtype=tl.float32) + for start in range(0, HW, BLOCK): + idx = start + offs + acc += tl.load(x_ptr + nc * HW + idx, mask=idx < HW, other=0.0).to(tl.float32) + tl.store(out_ptr + nc, (tl.sum(acc, axis=0) / HW).to(out_ptr.dtype.element_ty)) +``` +Typical SE-block region `mean → linear/conv1x1 → act → sigmoid → mul` can be split: +mean kernel (this) + elementwise tail kernel, two passes. + +## 7. Bias + activation epilogue (after a kept-in-aten linear/conv) + +Region = the ops AFTER `F.linear`/`torch.conv2d` (which stays in the graph): e.g. +`gelu(linear_out)` or `bn(conv_out) + relu-class`. Use template 1/3 with the activation: + +- exact GELU (`F.gelu`, approximate='none'): `0.5 * x * (1 + tl.math.erf(x * 0.7071067811865476))` +- tanh GELU (approximate='tanh'): `0.5 * x * (1 + tl.math.tanh(0.7978845608028654 * (x + 0.044715 * x * x * x)))` +- SiLU: `x / (1 + tl.exp(-x))` ; Sigmoid: `1 / (1 + tl.exp(-x))` +- LeakyReLU(s): `tl.where(x >= 0, x, s * x)` ; ReLU: `tl.maximum(x, 0.0)` +- These match eager within 1–2 ulp in fp32 — verify with the checker's smoke test; if a + strict fp32 variant fails, mirror eager more literally (e.g. use `tl.math.tanh`, keep + operation order identical). + +## 8. cat + elementwise (matches `torch.cat([a, b], dim) → ew-ops`) + +One kernel writes the output in segments; each segment loads from its source with that +source's strides: + +```python +@triton.jit +def _cat2_lastdim(a_ptr, b_ptr, out_ptr, R, NA, NB, BLOCK: tl.constexpr): + pid = tl.program_id(0) + N = NA + NB + offs = pid * BLOCK + tl.arange(0, BLOCK) + m = offs < R * N + r = offs // N + c = offs % N + from_a = c < NA + av = tl.load(a_ptr + r * NA + c, mask=m & from_a, other=0.0).to(tl.float32) + bv = tl.load(b_ptr + r * NB + (c - NA), mask=m & (~from_a), other=0.0).to(tl.float32) + v = tl.where(from_a, av, bv) + # ... fused elementwise ops on v here ... + tl.store(out_ptr + offs, v.to(out_ptr.dtype.element_ty), mask=m) +``` +For cat over dim 1 of NCHW use segment offsets with H*W blocks. ≥3 inputs: extend the +`tl.where` chain, or group segments with the same indexing/work into a small number of +launches writing slices of `out`. Do not assume one branch-heavy giant kernel is always +best for heterogeneous branches; bench it against grouped launches. Avoid one launch per +tiny segment unless the output work clearly amortizes the extra launch overhead. + +## 9. Embedding gather (matches positional `F.embedding(idx, weight, 0, None, 2.0, False, False)`) + +```python +@triton.jit +def _embedding(idx_ptr, w_ptr, out_ptr, T, D, BLOCK_D: tl.constexpr): + t = tl.program_id(0) + cols = tl.arange(0, BLOCK_D) + m = cols < D + row = tl.load(idx_ptr + t) # int64 index + v = tl.load(w_ptr + row * D + cols, mask=m, other=0.0) + tl.store(out_ptr + t * D + cols, v, mask=m) +``` + +## 10. Matmul + epilogue (`tl.dot`) — use SPARINGLY + +Only when the matmul is modest (K ≤ ~1024, M·N small enough that cuBLAS overhead matters) +or the epilogue chain is long; ALWAYS benchmark vs leaving matmul in aten. + +```python +@triton.jit +def _mm_epilogue(a_ptr, b_ptr, c_ptr, M, N, K, + sam, sak, sbk, sbn, + BM: tl.constexpr, BN: tl.constexpr, BK: tl.constexpr): + pm = tl.program_id(0) + pn = tl.program_id(1) + rm = pm * BM + tl.arange(0, BM) + rn = pn * BN + tl.arange(0, BN) + rk = tl.arange(0, BK) + acc = tl.zeros((BM, BN), dtype=tl.float32) + for k in range(0, K, BK): + a = tl.load(a_ptr + rm[:, None] * sam + (k + rk)[None, :] * sak, + mask=(rm[:, None] < M) & ((k + rk)[None, :] < K), other=0.0) + b = tl.load(b_ptr + (k + rk)[:, None] * sbk + rn[None, :] * sbn, + mask=((k + rk)[:, None] < K) & (rn[None, :] < N), other=0.0) + acc = tl.dot(a, b, acc) + # epilogue on acc here (bias add, activation, scale...) + tl.store(c_ptr + rm[:, None] * N + rn[None, :], + acc.to(c_ptr.dtype.element_ty), + mask=(rm[:, None] < M) & (rn[None, :] < N)) +``` +Note: `F.linear(x, w, b)` = `x @ w.T + b` — pass `w` with transposed strides +(`sbk = w.stride(1)...` i.e. read W[k-th col]) instead of materializing a transpose. +fp16/bf16 accumulate in fp32 (as cuBLAS does) — that's what `acc` does. `tl.dot` needs +block dims ≥16 and (for older Triton) M/N/K blocks multiples of 16 — mask handles edges. + +## 11. Two-kernel split for reduce→broadcast regions + +`y = f(x, g(reduce(x)))` where the reduction axis is large: kernel 1 computes the per-row +statistic into a small temp (`torch.empty((R,), ...)`), kernel 2 does the elementwise +combine. Both launched from ONE wrapper (still one pass / one route). + +## 12. In-wrapper scalar plumbing + +Pattern literals (eps, scales, dims) are fixed per pass — bake them as `tl.constexpr` +defaults or pass as kernel args from the wrapper. NEVER recover them from tensors. +If a region needs the same kernel with different constants per pass, parameterize via the +route string: `"bn_eps1e-05"` → parse in dispatch (string ops are fine in the wrapper). + +## 13. Reduced heavy ops — when a conv/matmul collapses to a cheaper primitive + +The vendor library is near-optimal for the *generic* op, not for the *reduced* form its +parameters may imply. Recognize the reduction from parameters and target the reduced primitive: + +- **Non-overlapping convolution** (`stride == kernel_size` per spatial dim, `kernel_size > 1`, + `padding == 0`): + windows are disjoint and reuse no input, so the conv is exactly a matmul of flattened + patches by the flattened weight: for output position `p` and out-channel `co`, + `out[p, co] = Σ_k patch[p, k] · W[co, k] + bias[co]`, where `k` ranges over + `in_ch · Π(kernel dims)` and `p` over `Π(output spatial dims)`. Gather each patch by index + arithmetic (NO `.contiguous()`), run `tl.dot`, add bias, then any fused tail + (flatten/transpose are absorbed for free by how you index/store; a contained + `cat`/`+pos`/`dropout(p=0)` tail can be folded into the epilogue). Skeleton (2-D windows; + drop/extend a dim for 1-D/3-D): + +```python +@triton.jit +def _patch_gemm(x_ptr, w_ptr, b_ptr, out_ptr, + P, CO, K, # patches, out-channels, contraction = in_ch*kH*kW + in_ch, H_in, W_in, kH, kW, H_out, W_out, + BM: tl.constexpr, BN: tl.constexpr, BK: tl.constexpr): + pm = tl.program_id(0); pn = tl.program_id(1) + rm = pm * BM + tl.arange(0, BM) # patch (output row) indices + rn = pn * BN + tl.arange(0, BN) # out-channel indices + acc = tl.zeros((BM, BN), dtype=tl.float32) + h_o = rm // W_out; w_o = rm % W_out # decode patch -> output spatial coords + base = (h_o * kH) * W_in + (w_o * kW) # top-left of each patch (stride==kernel), NCHW contiguous x + for k0 in range(0, K, BK): + rk = k0 + tl.arange(0, BK); km = rk < K + ic = rk // (kH * kW); r = rk % (kH * kW) + kh = r // kW; kw = r % kW + feat = ic * (H_in * W_in) + kh * W_in + kw + a = tl.load(x_ptr + base[:, None] + feat[None, :], + mask=(rm[:, None] < P) & km[None, :], other=0.0).to(tl.float32) + # W flattened to [CO, K] is contiguous: W[co, k] = w_ptr + co*K + k + b = tl.load(w_ptr + rn[None, :] * K + rk[:, None], + mask=km[:, None] & (rn[None, :] < CO), other=0.0).to(tl.float32) + acc += tl.dot(a, b) + acc += tl.load(b_ptr + rn, mask=rn < CO, other=0.0).to(tl.float32)[None, :] + tl.store(out_ptr + rm[:, None] * CO + rn[None, :], acc.to(out_ptr.dtype.element_ty), + mask=(rm[:, None] < P) & (rn[None, :] < CO)) +``` + +Scope: this specific patch-gather template applies to the **non-overlapping** conv only +(`stride == kernel_size`, `kernel_size > 1`, `padding == 0`); its index math does not model a +**1×1 conv** (already a pointwise matmul cuBLAS handles optimally — default to fusing its tail) +or a **depthwise/grouped/overlapping** conv (default to fusing the tail). Those defaults are +priors; a completed eval can overrule them for an off-regime shape. + +Correctness/perf notes (general): +- **`tl.dot(A, B)` computes `A @ B`, not `A @ Bᵀ`.** A conv weight is `[out_ch, in_ch, *kernel]`; + its flattened form is `[CO, K]`, so index it as `W[co, k]` with the contraction axis `k` on + the tile's row axis. A transpose slip here gives a large systematic `max_diff` on every + variant — always numeric-smoke-test (check_pattern `--smoke`) before an eval. +- Accumulate in fp32 (`.to(tl.float32)` after load), store in the eager output dtype. On fp32 + inputs the library may use TF32 tensor cores, so `tl.dot` will differ slightly — it passes + the fp32 *baseline* tolerance, not the strict bit-exact column; that is expected. +- Pick modest square-ish blocks (e.g. BM=BN=BK=64) as a default; `tl.dot` needs block dims + ≥16 with masked edges. If the sample has many variants, do NOT autotune (compile-cost × + variants risks the timeout) — pick one fixed config. +- If a contained `cat(cls_token)` shifts output rows by one (seq = 1 + P), treat row 0 as the + prepended token (load it, skip the GEMM) and map `patch = row - 1` for rows ≥ 1. +- Keep the conv's exact positional call form in `pattern` (it is C-bound: mirror stride/pad/ + dilation/groups positionally, per rule 2). diff --git a/skills/task-oriented/Passnet/skills/passnet-skill/SKILL.md b/skills/task-oriented/Passnet/skills/passnet-skill/SKILL.md new file mode 100644 index 000000000..512cfd34d --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-skill/SKILL.md @@ -0,0 +1,376 @@ +--- +name: passnet-skill +description: > + PassNet GPU kernel optimization via compiler passes. Design and implement Triton-based + optimization passes, create pass files under ./pass_dir/, self-evaluate with + pass_evaluator, and iterate to maximize GPU speedup. +--- + +You are an expert HPC Engineer specialized in Triton programming and GPU kernel optimization. + +Your task is to design and implement compiler optimization passes that achieve performance speedups on GPU. +You will analyze computation graphs, design pass structures to match target patterns, and implement +high-performance custom kernels using Triton. + +You are working in a specific problem directory where all your work is isolated. + +You are working on a PassNet (AI for Compiler) optimization task. + +**Goal:** +Optimize the target computation to achieve maximum performance speedup on GPU while maintaining correctness. + +**Key Task: Design an Ordered Sequence of Optimization Passes** +You have complete freedom to choose which operations to optimize and in what order. +Pass selection and ordering is a critical component - analyze the computation carefully to identify: +- Which operations can be fused or optimized independently +- What order maximizes performance gains +- How passes interact with each other + +**Your Working Directory:** +You are currently in the problem directory with the following structure: +- Pass files directory: ./pass_dir/ (NOTE: This directory is initially EMPTY. You need to CREATE the pass file from scratch) +- Evaluation script: ./entry.sh +- All file paths are relative to your current directory + +**Problem Statement:** +The problem statement is not pre-injected. You must read it from the current directory: +1. Read `graph_list.txt` to find the target graph path(s). +2. For each graph path, read `/model.py` (the computation to optimize) and `/weight_meta.py` (input tensor shapes, dtypes, and statistics). + +**General Approach:** + +1. **Analyze the Target Computation:** + - Study the graph information — it shows the exact computation to optimize + - model.py contains the computation pattern (e.g., Conv2D + ReLU, matmul + transpose) + - weight_meta.py contains input tensor shapes, dtypes, and statistics + - Use this information to understand what operations can be fused and optimized + +2. **Design the Optimization Pass(es):** + - Analyze the computation and identify independent optimization opportunities + - **IMPORTANT**: Create SEPARATE pass files for each independent optimization track + - For example, if the model has two independent operations, create two pass files: + * `FuseReduceSumDiv_dim2_keepdim.py` for normalization operations + * `FoldViewExpandToBroadcast_1_2_64_8_8.py` for view/expand operations + - Use descriptive names that indicate what the pass optimizes + - Each pass file should have three functions: + * `pattern`: A function that matches ONE specific computation pattern + * `replacement_args`: A function that extracts necessary arguments from matched nodes + * `replacement_func`: Returns a custom implementation that's faster than the original + +3. **Create the Pass Configuration File:** + - **CRITICAL**: You MUST create `./pass_dir/sorted_output_pass_rule_names.json` + - This defines your optimization strategy - which passes to apply and in what order + - The order matters! Passes are applied sequentially, so consider dependencies and performance impact + - Format: JSON array with pass names EXACTLY matching your Python filenames (without .py) + - Example: If you create these Python files: + * `./pass_dir/FuseReduceSumDiv_dim2_keepdim.py` + * `./pass_dir/FoldViewExpandToBroadcast_1_2_64_8_8.py` + Then create `./pass_dir/sorted_output_pass_rule_names.json`: + ```json + ["FuseReduceSumDiv_dim2_keepdim", "FoldViewExpandToBroadcast_1_2_64_8_8"] + ``` + - **The evaluation framework requires this file to discover and load your passes** + +4. **Implement the Optimized Kernel:** + - Write a high-performance kernel using Triton (or other GPU programming frameworks) + - Consider tensor shapes from weight_meta.py when choosing tile/block sizes + - Optimize for memory coalescing, shared memory usage, and GPU occupancy + - Ensure semantic equivalence - the kernel must produce the same results as the pattern + +5. **Test and Iterate:** + - Use pass_evaluator to run evaluation (no arguments needed) + - Check three metrics: pass matching, correctness, and speedup + - Adjust your implementation based on results + - Try different optimization strategies and kernel configurations + - Continue iterating to maximize speedup + +**Technical Requirements:** +- You MUST create at least one pass file under ./pass_dir/ and it must be importable by the evaluation framework (no syntax errors, missing imports, or unresolved symbols). +- You MUST create ./pass_dir/sorted_output_pass_rule_names.json. + - It must be a JSON array of strings. + - Each string MUST exactly equal a pass Python filename without .py. + - Every pass you want applied MUST appear in this list, in the exact order you want them executed. +- Pattern outputs MUST include every value that is observable outside the matched subgraph (e.g., any intermediate that appears in the model's return). +- replacement_func() MUST be a zero-argument function that returns a callable function object (DO NOT call it). +- API Validation (enforced on all functions except `pattern()` and `replacement_args()`): + - Allowed: tensor allocation APIs only — `torch.empty`, `torch.empty_like`, `torch.zeros`, `torch.zeros_like`, `torch.ones`, `torch.ones_like`, `torch.full`, `torch.full_like`, `torch.as_tensor`. + - Blocked: all other `torch.*` calls and imports of `torch.nn`, `torch.ops`, `torch.autograd`. Use Triton kernels for computation. + +**Creating Pass Files:** +Write pass files directly into `pass_dir/` using your built-in file editing capabilities. +Create each `.py` pass file and `sorted_output_pass_rule_names.json` under `pass_dir/`. +The `draft/` directory is available for exploratory work and evaluation response logs. + +**Pass File Structure:** +Your pass file must follow this structure for the framework to work correctly: +```python +import torch +import triton +import triton.language as tl + +# Pattern matching function +def pattern(arg1, arg2, ...): + """ + Define the computation pattern to match + SPECIAL NOTE: The operations in this function MUST mirror the operations in model.py exactly (including positional vs keyword arguments, op variants, and dataflow). + e.g.: + if model.py has `tmp1 = torch.conv2d(input_tensor, weight_tensor, bias_tensor, (1, 1), (0, 0), (1, 1), 1)`, + pattern MUST also use positional arguments for stride, padding, dilation, and groups, not keyword arguments. + **Wrong case**: result = torch.conv2d(input_tensor, weight_tensor, bias_tensor, stride=(1, 1), padding=(0, 0), dilation=(1, 1), groups=1) + **Right case**: result = torch.conv2d(input_tensor, weight_tensor, bias_tensor, (1, 1), (0, 0), (1, 1), 1) + """ + result = ... # operations to match + return result + +# Argument extraction function +def replacement_args(arg1, arg2, ...): + # Extract and return arguments needed for the replacement + return (arg1, arg2, ...) + +# Your optimized kernel +@triton.jit +def optimized_kernel(...): + # High-performance implementation + ... + +# Kernel wrapper (MUST be decorated with @torch.fx.wrap) +@torch.fx.wrap +def kernel_wrapper(...): + # Set up grid and launch kernel + optimized_kernel[grid](...) + return result + +# Replacement function (NO arguments, returns function reference) +def replacement_func(): + return kernel_wrapper # Return the function, not a call +``` + +There is a reference optimization passes for Triton kernel. +Give unoptimized pass: + ```python + import torch + + def pattern(x, y): + return x+y + + def replacement_args(x, y): + return (x, y) + + def replacement_func(): + pass + ``` + + Output optimization Pass: + ```python + + @triton.jit + def triton_add_kernel( + x_ptr, + y_ptr, + out_ptr, + n_elements, + BLOCK_SIZE: tl.constexpr, + ): + # Each program handles a contiguous block of data of size BLOCK_SIZE + block_start = tl.program_id(0) * BLOCK_SIZE + offsets = block_start + tl.arange(0, BLOCK_SIZE) + mask = offsets < n_elements # Mask to ensure we don't go out of bounds + # Load + x = tl.load(x_ptr + offsets, mask=mask, other=0.0) + y = tl.load(y_ptr + offsets, mask=mask, other=0.0) + # Calculate + out = x + y + # Store + tl.store(out_ptr + offsets, out, mask=mask) + + @torch.fx.wrap + def triton_add(x, y): + N = x.numel() + BLOCK_SIZE = 1024 + num_programs = (N + BLOCK_SIZE - 1) // BLOCK_SIZE + + out = torch.empty_like(x) + + triton_add_kernel[(num_programs,)]( + x_ptr=x, + y_ptr=y, + out_ptr=out, + n_elements=N, + BLOCK_SIZE=BLOCK_SIZE, + ) + + return out + + def replacement_args(x, y): + return (x, y) + + def replacement_func(): + return triton_add + ``` + +**Pattern Matching Guidelines:** + +Pattern matching is performed over the exact dataflow structure of the computation graph. +Any intermediate value that is observable outside the matched subgraph—in particular, values that appear in the model's return—must be explicitly produced by the pattern. + +**IMPORTANT**: Do NOT include cleanup statements like `tmp_x = None` in your pattern. + +Example: Given a model: +```python +class Model(torch.nn.Module): + def forward(self, in_0): + ... + tmp_5 = tmp_4.transpose(-1, -2) + tmp_9 = tmp_5 @ tmp_6 + return (tmp_5, tmp_8, tmp_9) +``` + +You decide to optimize `transpose + matmul` pattern. The correct pattern is: +```python +def pattern(a, b): + t = a.transpose(-1, -2) + out = t @ b + return t, out + +def replacement_args(a, b): + return (a, b) + +def replacement_func(): + pass +``` + +❌ WRONG - fuses operations without creating observable intermediate `t`: +```python +def pattern(a, b): + out = a.transpose(-1, -2) @ b + return out + +def replacement_args(a, b): + return (a, b) + +def replacement_func(): + pass +``` + +**Best Practices:** +- Create separate pass files for independent optimization opportunities (don't try to optimize everything in one pass) +- Pattern matching is strict - only include actual operations, exclude `tmp_x = None` cleanup statements +- Each pattern should match ONE specific operation or fusion opportunity +- Each pattern should include at least one Triton kernel implementation +- Pay attention to what the model returns - your pattern must return the same structure +- ALWAYS create sorted_output_pass_rule_names.json listing all your pass files +- When iterating, focus on optimizing the kernel implementation rather than changing pattern/replacement_args +- Test early and often with pass_evaluator to catch issues +- Learn from evaluation feedback - if pattern doesn't match, check your pattern function carefully +- For correctness failures, verify your kernel logic and data types +- For speedup optimization, first analyze the performance bottlenecks of the Triton kernel, then progressively apply optimizations such as autotuning configurations, re-tile for better parallelism (e.g. change grid dimensions or size, the kernel should be modified accordingly.), and kernel fusion. +- When the pattern matches, you should focus on optimizing kernel performance, such as adding @autotune configs to Triton functions or tuning the parameters in those configs. +- (Optional) If some Passes fail to match while others sharing similar logic succeed, consider consolidating them: use a parameterized pattern for Passes that differ only in a scalar constant (e.g., a division scale), or extract a shared Triton kernel into a separate file under pass_dir/ and import via `from pass_dir.your_kernel_file import your_func`. +- If the above consolidation still fails due to replacement_func_limit dropping your passes: Use the shared replacement_func routing technique — make ALL pass files share the SAME `replacement_func()` returning a single `@torch.fx.wrap` dispatch wrapper, and differentiate each pass by appending a **route string** as the last argument in `replacement_args()` (e.g., `return (x, "route_a")`). Inside the shared dispatch wrapper, use `if/elif` on the route string to call the corresponding Triton kernel. Every pass file must define the full dispatch wrapper with all route branches (the elif branches for other routes can call placeholder private functions — they never execute in that pass's context). This way `replacement_func()` is identical across all passes, so `output_pass_replacement_func_limit` never drops any of them. + +## Benchmark API + +You are running in a **sandboxed** develop directory created by the PassNet Benchmark API. +GPU evaluation is available exclusively through the Benchmark API at `$API_URL`. + +### Environment + +- **Sandbox**: You are inside a Bubblewrap mount namespace. The current directory is your + isolated develop directory. You can only see this sample's resources — other samples are + invisible. +- **Writable**: `pass_dir/` (pass files ready for submission), `draft/` (exploratory work, + evaluation response logs). +- **Read-only**: `graph_list.txt`, graph directories (`model.py`, `weight_meta.py`), + `pass_bench/`, `entry.sh`, `evaluations/`. +- **No GPU access**: GPU evaluation runs on the API server, not locally. +- **Feedback scripts** (if available): `~/.claude/skills/passnet-feedback/scripts/` + (`analyze_graph.py`, `check_pattern.py`, `parse_eval_log.py`). They auto-detect the + PassNet root from the `entry.sh` symlink. + +### Check Service + +`GET $API_URL/health` checks whether the Benchmark API is available. + +```bash +curl -s "$API_URL/health" | python3 -m json.tool +``` + +Proceed only when the response reports `"status": "ok"`. + +### Read the Problem + +Read the problem directly from the develop directory — resources are symlinked in: + +```bash +cat graph_list.txt +cat /model.py +cat /weight_meta.py +``` + +### Evaluate + +`POST $API_URL/evaluate` validates and evaluates the pass files in `pass_dir/`. + +Construct the payload from the current contents of `pass_dir/`: + +```bash +python3 -c ' +import json, os +from pathlib import Path + +root = Path("pass_dir") +files = { + path.name: path.read_text(encoding="utf-8") + for path in sorted(root.iterdir()) + if path.is_file() and path.suffix in {".py", ".json"} +} +print(json.dumps({ + "develop_id": os.environ["DEVELOP_ID"], + "subdirectory": os.environ["SUBDIRECTORY"], + "sample_path": os.environ["SAMPLE_PATH"], + "files": files, +})) +' | curl -s --max-time 1800 -X POST "$API_URL/evaluate" \ + -H 'Content-Type: application/json' \ + -d @- \ + | tee draft/last_evaluation_response.json | python3 -m json.tool +``` + +**Request fields:** +- `develop_id`: the assigned `$DEVELOP_ID`. +- `subdirectory`: the assigned `$SUBDIRECTORY`. +- `sample_path`: the assigned `$SAMPLE_PATH`. +- `files`: object mapping each top-level `pass_dir/` filename to its UTF-8 text content. + Include all `.py` and `.json` files. + +**Response fields:** +- `returncode`: evaluation process return code (0 = success). +- `pass_matched`: whether any submitted pass matched the graph. +- `aggregated_score`: score object with `id` and `score` fields. +- `result_dir`: path relative to the develop directory containing evaluation artifacts + (`validation.log`, `aggregated_score.json`, `submission/` copy). + +**After each evaluation**, results are persisted to `evaluations/_/`. +Inspect `validation.log` and `aggregated_score.json` there; use +`~/.claude/skills/passnet-feedback/scripts/parse_eval_log.py` to parse them. + +### Evaluation Budget + +`MAX_EVALUATIONS` is the hard limit on `POST /evaluate` requests for this task. +**Count every request** you send to `/evaluate`, including requests that fail validation +or evaluation. Never send more than `MAX_EVALUATIONS` requests. Use the remaining +attempts deliberately and leave the best final implementation in `pass_dir/`. + +### Finish + +When you have used all `MAX_EVALUATIONS` attempts, or when you decide to stop optimizing, +review all results in `evaluations/`, select the version with the highest score, and copy +the files from that version's `submission/` directory into `pass_dir/` as the final +submission. + +In your final reply, report: +- Final score +- Whether the pass matched +- Pass files created +- Optimization strategy +- Failure reason, if applicable diff --git a/skills/task-oriented/Passnet/skills/passnet-solve/SKILL.md b/skills/task-oriented/Passnet/skills/passnet-solve/SKILL.md new file mode 100644 index 000000000..472329b1d --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-solve/SKILL.md @@ -0,0 +1,288 @@ +--- +name: passnet-solve +description: > + END-TO-END playbook for solving one PassNet sample: analyze the computation graphs, + decide the optimization strategy (what to fuse, what to replace, what to leave alone), + drive the iteration loop, and maximize the sample score. This is the ENTRY skill — + invoke it first for any PassNet optimization task; it tells you when to use + passnet-pattern-fusion, passnet-triton-opt and passnet-feedback. +--- + +You are optimizing one PassNet sample. A sample = several *graph variants* of the same +subgraph (different dtypes float32/float16/bfloat16 and input seeds/batch dirs). You write +compiler passes into `pass_dir/`; the evaluator pattern-matches them into each variant's FX +graph, checks correctness vs eager, and measures speedup. + +## 0. How the score works (drives every decision) + +Per graph variant, the *rectified speedup* is: + +| outcome | value | +|---------------------------------------------|------------------| +| no pass matched, or crash/runtime error | **0.1** | +| matched but wrong numerics (baseline tol) | **≈0.15** effective | +| matched + correct, e2e speedup `s` | **≈ s** (even if s < 1) | + +Sample score = geometric mean across ALL variants (then a tolerance-weighted aggregate that +≈ the geomean). Read `references/passbench-internals.md` for the exact math, the harness +mechanics, anti-cheat rules, and evaluation modes. + +Consequences — internalize these: +1. **A matched, correct pass with speedup 0.8 scores 8× better than no match (0.1).** + Never finish with an empty / non-matching `pass_dir`. +2. One crashing variant (0.1) drags the geomean of the whole sample hard. Consistency + across ALL variants beats a hero kernel on one variant. +3. Speedup is **end-to-end wall clock** (Python overhead included), median of 100 trials, + vs eager. The compiled path pays a fixed per-call tax (dynamo guards + FX + interpretation + wrapper + Triton launch). **The tax scales with graph size and is + machine-dependent: ~40–70 µs on a 3-op graph, ~150–230 µs measured on a 13-node graph. + Calibrate it from your first eval** (`tax ≈ compiled_e2e − your_kernel_µs − + remaining_eager_ops_µs`) and reuse it in ceiling estimates. You only win if the aten + work you absorb into your kernel(s) exceeds the tax. +4. **A floor pass is downside insurance, not proof of upside.** Use it to avoid the 0.1 + cliff and calibrate tax. If a completed eval shows a floor or narrow pass is matched + and correct but slower than eager, do not blindly tune the same small kernel. Either + attempt a clearly larger legal region that should amortize the fixed tax, or record an + overhead-bound stop when no such region exists. + +## 1. Mandatory workflow + +### Step 1 — Detect the environment mode (once, at start) + +**Explicit task instructions override detection** — if your prompt says which mode to use +(or gives SVC/SAMPLE variables), use that mode; a service answering on the port may +belong to another tenant/GPU on a shared box. + +```bash +curl -sS --max-time 5 http://127.0.0.1:${PASSNET_API_PORT:-8968}/health +``` +- Response contains `"mode": "sample_access_service"` → **service mode**: all file I/O via + HTTP with `?sample_path=$SAMPLE`; GPU eval via `POST /evaluate`. See + `references/passbench-internals.md` §"Evaluation modes" for the exact curl recipes. +- Response is `{"status":"ok"...}` without that mode → **API-server mode** (same endpoints, + no `sample_path` param). +- No response (or you were told to work locally) and you are in a sample directory (has + `entry.sh`, `graph_list.txt`) → **local mode**: edit `pass_dir/` directly; evaluate with + `bash entry.sh` (or `pass_evaluator` if available). Local mode always works from inside + a sample dir regardless of what services are running. + +In service mode, still author files in a local scratch dir first, run the pre-flight +checker on them, then POST them. The service is the source of truth: after upload, call +`/files` and confirm the expected JSON manifest and at least one pass file are visible +before `/evaluate`. Use bounded curl timeouts. In Codex managed sandboxes, try normal +localhost service `curl` first and continue normally if it succeeds. If it fails with +HTTP000, curl exit code 7, an empty body, or a connection-style error, retry once and +then use the approved escalation path if available. A sandboxed connection error does +not prove the service is down. Keep `curl` as its own command segment, capture response +bodies to files, and capture HTTP status separately so status text never corrupts a JSON +body. + +### Step 2 — Read the problem + +Get `graph_list.txt` and every variant's `model.py` + `weight_meta.py` (locally or via +`/problem`). Note for each variant: ops and their **exact textual call form**, tensor +shapes, dtypes. Variants usually share the op structure; shapes may differ — check, it +decides whether shape literals may appear in patterns. + +### Step 3 — Analyze before writing anything + +Run the analyzer from passnet-feedback (CPU-only works; GPU enables timings): + +```bash +SCRIPTS=/.claude/skills/passnet-feedback/scripts +python3 $SCRIPTS/analyze_graph.py --sample-dir [--bench] +``` + +It prints, per variant, the dynamo graph with per-node **matchability** classification and +(with `--bench`, GPU) per-node eager timings. This tells you (a) which nodes a pattern can +legally bind to, (b) where the time actually goes — the bottleneck. + +### Step 4 — Choose the strategy (decision policy) + +Work through this in order: + +0. **Heavy-op regime check — do this before labeling any dominant conv/matmul "leave in the + vendor lib."** A vendor kernel is near-optimal only inside the shapes/parameters it tuned for + and for the *general* form of its op; you can beat it only when the concrete parameters put + the op OUTSIDE that regime AND equal a cheaper computation. The clearest checkable win to + actively look for: a **convolution with `stride == kernel_size` AND `kernel_size > 1` AND + `padding == 0`** (non-overlapping windows) — disjoint receptive fields make it exactly a dense + patch matmul while cuDNN runs a heavier general path, so a `tl.dot` patch-gather kernel can win + big; treat it as an ambitious anchor AFTER shipping the tail/floor. For ops in the vendor's + sweet spot — 1×1 conv (already a cuBLAS GEMM), overlapping conv (im2col would blow up memory), + depthwise/grouped, general dense matmul — the reliable default is leave-in-aten + fuse-tail; + rewriting them usually loses. These are strong priors, not absolute bans: the completed eval + is the arbiter, so if parameters give a concrete reason to expect a win you may try — but only + after a floor is banked and only keep it if the eval beats the floor. Don't conclude a + conv/matmul is "unmatchable" (they are C-bound and match fine); see passnet-orchestrate §2.2. + +1. **Partition the graph** into maximal *matchable regions*: connected runs of matchable + nodes (see passnet-pattern-fusion for the matchability rules) where the region has + exactly ONE output consumed outside it. For normal callable patterns, kwargs-form + Python-level F.* calls (for example `F.relu(x, inplace=False)`) usually become region + boundaries because the pattern tracer normalizes them differently. If such a node blocks + a high-value single-output region, use passnet-pattern-fusion's manual FX `GraphModule` + pattern escape hatch and require real `check_pattern.py` proof; otherwise leave it in + eager and fuse around it. Do not absorb `inplace=True` nodes merely because manual FX can + match them; either reproduce the eager mutation exactly across repeated calls or leave the + in-place node outside and fuse the downstream side-effect-free region. +2. **Estimate the win for each region** (numbers from `--bench`, or the heuristics in + passnet-triton-opt §"performance model"): + `gain ≈ (sum of eager µs of absorbed nodes) − (fixed tax ~50 µs amortized over regions) − (your kernel µs)`. + - Many small elementwise/norm ops in one region → fuse them into ONE Triton kernel: + best case, this is where big speedups live. + - Long memory-bound elementwise or broadcast-affine-activation chains are now priority + upside targets when they can be expressed as one legal single-output region. A floor + still comes first for downside protection, but do not stop at a narrow floor on these + shapes: attempt the larger legal region that absorbs the broadcast/layout producers + and multiple affine, activation, residual, or normalization consumers unless + containment, matchability, or numeric pre-flight proves it illegal. + - Normalization-affine compute prefixes and concat→normalization-affine prefixes are also + priority likely-amortizable targets when the prefix has one legal output and enough + output work to pay the launch/framework tax. Keep a normalization-only floor first, then + evaluate the broader prefix; repeated results show this family can win strongly, but + small or branchy instances can still be correct-but-slow or regress and must be reverted + to the best completed state. + - `conv` / `matmul`: cuDNN/cuBLAS are near-optimal **only inside the regime they tuned for + and for the op's general form**. Reliable default: leave the heavy op in aten and fuse its + elementwise *tail* (bias, norm, activation, residual add) into one kernel; re-implementing + the op competes with the library on its home turf and usually loses. + - **When a rewrite CAN win — the clearest case: a convolution whose windows do not overlap** — + `stride == kernel_size` along each spatial dim AND `kernel_size > 1` AND `padding == 0`. + Then the receptive fields are disjoint, so the conv is exactly a matmul of flattened patches + by the flattened weight (`patches @ weight.reshape(out_ch, -1).T`, plus bias). cuDNN still + runs a general im2col/implicit-GEMM path with layout overhead it cannot skip, so a `tl.dot` + patch-gather kernel that also folds the flatten/transpose/tail into its epilogue can win by + a wide margin (the patch-embedding stem pattern). This is a real algorithm change, not a + reimplementation of the same work. Treat it as an ambitious region AFTER the floor. + - **Why the same rewrite usually loses elsewhere (so the default holds):** + * **1×1 conv** (kernel_size 1) already IS a pointwise matmul cuBLAS runs directly — no + general-conv overhead to remove, so `tl.dot` competes head-to-head and rarely wins. + * **overlapping conv** (stride < kernel_size): a matmul form needs im2col, materializing a + ~kernel_volume× larger matrix — the blow-up cuDNN's implicit-GEMM/Winograd avoids. + * **depthwise / grouped conv** (groups > 1): memory-bound with specialized vendor kernels; + grouped indexing is bug-prone. + * **general dense matmul**: cuBLAS's sweet spot. + - **These are strong priors, not bans — the completed eval is the arbiter.** If a specific + op's parameters give you a concrete reason to expect it sits outside the vendor's efficient + regime (a tiny/degenerate matmul, a memory-bound op fusable with a long tail into one + kernel), you MAY try the rewrite — but only after a floor is banked, and keep it only if a + completed eval beats that floor. Don't burn budget rewriting a sweet-spot op on a hunch; + don't skip the non-overlapping case just because nothing forbids it. + - Do NOT assume a heavy op is unmatchable: `conv1d/2d/3d`, `matmul`, `F.linear` are C-bound + functions the pattern tracer binds to fine — mirror the exact positional call form and they + match (see passnet-pattern-fusion §3). "It's a big cuDNN op" is never a reason to conclude + the pattern can't match. + Discipline is unchanged: ship the safe tail/floor first, then attempt any rewrite as an + additional ambitious region and REVERT (Step 7) if a completed eval does not beat the floor. + - Pure layout ops (`view`/`reshape`/`permute`/`transpose`/`unsqueeze`) cost ~0 in eager; + absorbing them only pays when it lets you bridge two compute ops into one kernel, or + erase a `.contiguous()` materialization. + - Layout/view/reshape-only regions and split/fanout tails are usually downside caps: + useful when they are the only safe way to match, but not a reason to spend tuning + budget unless they bridge real compute or remove an actual materialization. + - Be careful with vendor-heavy tails and layout materialization: a larger-looking region + can lose when it replaces a vendor-tuned producer or forces eager-metadata layout work + to become a real copy. If the completed larger-region eval is correct but slower, revert + to the best completed state and record the caveat instead of tuning blindly. +3. **Total-eager-time sanity check**: if the variant's whole eager forward is tiny + (< ~150 µs e2e, i.e. a couple of cheap ops) the fixed tax means the ceiling is < 1. + Don't burn time chasing >1; ship the safest correct pass (floor pass) and move on — + 0.85 is a fine score for such a sample. + If a narrow or floor pass has a completed eval, matches everywhere, is correct, and + is still slower than eager, classify it as overhead-bound after checking region + availability; keep it as the downside cap, revert any lower-scoring wider attempt, + and stop unless there is a clearly larger legal region that absorbs enough eager work + per launch to plausibly beat the tax. +4. **Skip-list**: never include in a pattern: RNG ops (`torch.rand`, dropout with + `training=True`), data-dependent control flow, ops you cannot reproduce faithfully + at the dtype's baseline tolerance (see tolerance table in references). Dropout with + `training=False` IS safe (identity) when its call form is matchable. +5. **Multi-variant check**: a pattern containing shape literals (for example, a view with fixed dimensions) + only matches variants with those exact literals. If shapes differ across variants, + either avoid literal-bearing nodes in patterns, or write one pass per shape family — + every variant must end with ≥1 matching pass. + +### Step 5 — Ship the floor pass FIRST + +Housekeeping when starting on a sample: `rm -rf pass_dir/__pycache__` and remove any +leftover `.py`/`.json` from previous occupants — `pass_dir` is imported as a package and +stale files/pyc can shadow your modules or load as unexpected extra passes. + +Before any ambitious work, write + verify + evaluate the simplest possible pass: +**one matchable compute node** (prefer `batch_norm`/`layer_norm`/`linear`-tail/`gelu`/ +`sigmoid`-class single op; a straightforward Triton kernel you cannot get wrong). +Once it scores (anything ≥ ~0.7), the sample's downside is capped. Keep it in `pass_dir` +until something better replaces it; ALWAYS re-evaluate after replacing. + +### Step 6 — Build the upside, iterating + +- Write fusion passes per region → use **passnet-pattern-fusion** (file contract, the 10 + hard rules, shared-dispatch architecture for multiple passes, kernel templates). +- ALWAYS pre-flight before GPU eval (`check_pattern.py`, from **passnet-feedback**): + it verifies load → AST validation → match-per-variant → replacement wiring → a one-shot + numeric smoke test, in seconds. Only then run the real evaluation. +- Tune slow-but-correct kernels with **passnet-triton-opt** only when the region still has + plausible headroom. A completed eval that is correct but slow on a maximal tiny or + layout-heavy region is usually a stop/revert decision, not an invitation to sweep block + sizes. +- Interpret each evaluation with **passnet-feedback** (`parse_eval_log.py` gives per-variant + status/speedups + estimated sample score + failure classification). Only a nonempty, + JSON-parseable evaluation response counts as completed; if the response is empty, + interrupted, or non-JSON, leave metrics unknown and fix service/upload/access state first. + +### Step 7 — Iterate with discipline + +- One change-class per iteration (match fix | numeric fix | perf tune); re-check, re-eval. +- Track the best-scoring `pass_dir` state; if an "improvement" lowers the score, revert. + In service mode keep local copies of every uploaded version so you can restore. +- If a larger legal region completed evaluation but scored lower than the floor or narrow + state, restore the best completed-evaluation state immediately and do not stack more + passes on the worse state. +- If a broad manual-FX region is matched and correct but unstable or slower, try the narrower + side-effect-free region before tuning blindly, especially when the broad region absorbed an + in-place node or added branch-heavy index logic. +- Typical sane budget: ≤ 6–8 GPU evaluations per sample. Each full eval costs minutes + (100 trials × variants); the 600 s timeout is real — samples with many variants + (some have 100+ graphs) can time out if your kernel compiles many autotune configs. + For such samples: NO autotune, one fixed config, minimal passes. +- Stop when: score ≥ your estimated ceiling × ~0.9, budget is exhausted with a correct + matching pass in place, or completed evals show a correct narrow/floor pass is + overhead-bound and no larger legal fusion region remains. + +## 2. Strategy cheat-sheet by graph shape + +| graph looks like | strategy | +|---|---| +| chain of elementwise + norm ops (MLP tail, norm+act+dropout) | one fused elementwise/row kernel over the maximal matchable region | +| broadcast/layout producer + affine/activation/residual consumers writing a large output | priority target: one legal single-output fused memory-bound kernel; verify it does not force a returned view or metadata-only layout into a costly materialization | +| concat or multi-branch producer → inference normalization → affine/activation | priority target when containment gives one normalized/final output; leave in-place downstream activations outside, and only include kwargs-form Python functional activations when a manual FX pattern pre-flight proves the exact region matches | +| `linear → (gelu/relu/silu) [→ dropout]` | fuse linear's tail if linear is big (keep cuBLAS), or `tl.dot`+epilogue if small; both: measure | +| `conv` with `stride == kernel_size`, `kernel_size > 1`, `padding == 0` (non-overlapping windows) | reduces to a patch matmul: gather disjoint patches + `tl.dot` against the flattened weight, and fuse the flatten/transpose/tail into the same kernel. Treat the conv as the region ANCHOR, not a wall. Ship a tail floor first | +| `conv → bn/act/add` (1×1, depthwise/grouped, or overlapping conv) | leave the conv in aten; fuse the tail ops into one kernel. Rewriting the conv usually loses (vendor sweet spot); the non-overlapping `stride==kernel_size`, `kernel>1` conv above is the case that IS worth rewriting | +| attention-ish (`matmul → softmax(method form) → matmul`) | softmax row-kernel (method-form softmax matches!); large generic matmuls stay aten; consider fusing scale/mask into the softmax kernel | +| `cat`/`stack` of tensors + elementwise or repeated producers | write the final output directly; compare one segmented kernel with a small number of grouped per-branch launches when branches have very different indexing/work | +| mean/sum/norm reductions + broadcast arithmetic (L2-norm, RMSNorm, sq-distance) | single row-reduction kernel producing the final output (verified ~1.6–2.5× on ms-scale tensors) | +| `expand`/broadcast → binary op writing a LARGE output | strong target even alone: aten broadcast-elementwise kernels often run 1.5–2.5× off the bandwidth roofline; a shape-specialized Triton kernel + erased layout nodes wins big (verified 2.3×) | +| mostly `view/permute/contiguous` + 1 compute op | replace the compute op only; or compute-op kernel that reads with source strides (absorbing a `contiguous`) | +| single huge op that is a generic dense matmul / 1×1 / depthwise / overlapping conv | floor pass on an adjacent/auxiliary node; accept ≈1.0; reimplementing usually loses. (A non-overlapping `stride==kernel_size`, `kernel>1` conv is the case worth rewriting — see that row.) | +| graph contains `torch.rand` / training-mode dropout | keep RNG nodes OUT of patterns; optimize around them | + +## 3. Red lines (read references for the full list) + +- Real computation must happen in **Triton kernels**. The wrapper may only use the + whitelisted `torch.empty/zeros/ones/full[_like]/as_tensor` allocators plus shape/stride/ + dtype/device accessors — every other aten op on the (poisoned) inputs raises during the + anti-cheat warmup, failing the variant. +- Do NOT evade validation (no `getattr(torch, ...)` laundering, no `no_dispatch()`, no + monkeypatching the harness). It's classified as hacking behavior; assume it will score 0 + on the leaderboard even if it slips through locally — and it's usually slower anyway. +- `replacement_func()` must return a stable module-level function, and with + `output_pass_replacement_func_limit: 1` **all your passes must return the SAME function + object** — the shared-dispatch architecture in passnet-pattern-fusion is mandatory the + moment you have ≥2 pass files. + +## 4. Final report + +End with: final score (last eval), pass_matched, per-variant speedup summary, list of pass +files, 1-paragraph strategy description, and (if below ceiling) what blocked further gains. diff --git a/skills/task-oriented/Passnet/skills/passnet-solve/references/passbench-internals.md b/skills/task-oriented/Passnet/skills/passnet-solve/references/passbench-internals.md new file mode 100644 index 000000000..d5c367752 --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-solve/references/passbench-internals.md @@ -0,0 +1,188 @@ +# PassBench harness internals — verified mechanics + +Everything below was verified against `pass_bench/` source and live runs. File pointers use +repo-relative paths. + +## 1. Evaluation pipeline (`entry.sh` → `pass_bench.torch.test_compiler`) + +Per graph variant (each runs in a FRESH subprocess): +1. `model.py` is loaded; inputs are replayed from `weight_meta.py` (see §6). +2. `torch.compile(model, backend=PassMgrBackend)` traces the model with dynamo → an FX + graph in **torch IR** (the ops exactly as written in `model.py` source — not aten IR). +3. `PassMgrBackend` loads pass files from `pass_dir/` in the order given by + `sorted_output_pass_rule_names.json`, applies each as a pattern-replacement pass. + If NO pass modifies the graph → RuntimeError → that variant scores 0.1 and the whole + sample early-exits ("Pass testing early exits on pass mismatch"). +4. **Warmup 25 calls with anti-cheat poison ON** (§4), untimed. +5. **100 timed trials with poison OFF**: per-trial `e2e` (host wall clock) and `gpu` + (cuda events) times; medians reported. Eager model measured the same way. + Note: each call is preceded by `torch.manual_seed(1024)` — RNG ops are reproducible + but consume the RNG stream; do not change how many RNG ops execute. +6. Correctness: `torch.allclose` sweep over tolerance levels + dtype equality of outputs + (`[Datatype]` lines) + `torch.equal`. `[Result] status: success|failed`. +7. `aggregate_es_scores.py` computes the sample score from the log. + +Config used by `entry.sh`: +`output_pass_pattern_limit: 100`, **`output_pass_replacement_func_limit: 1`** (critical, §3), +warmup 25, trials 100, `--device cuda`. + +## 2. Scoring math (`pass_bench/analysis_util.py::calculate_scores` + `aggregate_es_scores.py`) + +- Speedup used: **`[Speedup][e2e]` = eager_e2e_median / compiled_e2e_median**. +- Tolerance level `t` maps to `(rtol, atol)` per dtype + (`datatype_tolerance_config.py::get_precision`): `rtol = 10^(t·k)` with k = + {fp16: 0.6, bf16: 0.3592, fp32: 1.1772}; `atol = 10^t`. The **baseline** is `t = -5`: + | dtype | rtol | atol | + |---|---|---| + | float32 | 1.3e-6 | 1e-5 | + | float16 | 1.0e-3 | 1e-5 | + | bfloat16 | 1.6e-2 | 1e-5 | +- Per variant per `t`: correct → rectified speedup `s` (p=0 ⇒ no slowdown squaring); + failed → `0.1`; at `t ≥ 1` failures get "tolerated" (accuracy at t≥1, runtime/compile at + t≥3 → value 1). +- ES(t) = geometric mean of rectified speedups over all variants in the sample. +- Final = `10^(Σ w_t·log10 ES(t))` with normalized weights + `{-10..-6: 0.001 each, -5: 1, -4: 1, -3: 1, -2: 0.8, -1: 0.64, 0: 0.512, 1: 0.4096, + 2: 0.32768, 3: 0.262144, 4: 0.001}`. +- Net effect: **correct at t=-5 with speedup s ⇒ sample ≈ 0.998·s** (strict t≤-6 levels are + ~weightless). Matched-but-wrong ⇒ ≈ 0.147. No match / crash ⇒ 0.1. + +## 3. PassMgrBackend mechanics (`pass_bench/torch/backend/pass_mgr_backend.py`) + +- Pattern and replacement are traced with `ForceArgsTracer` + (`custom_replacement.py`): for `call_function` nodes whose target has an + `inspect.signature` that binds, args are rewritten to FULL POSITIONAL with defaults + filled, kwargs emptied. `call_method` nodes and C-bound functions (no signature) keep + exactly the written form. The TARGET graph (from dynamo) is NOT normalized → the + matchability taxonomy in passnet-pattern-fusion. +- Matching: `SubgraphMatcher` with `match_output=False, match_placeholder=False, + remove_overlapping_matches=True, ignore_literals=False`. + - Literals must be EXACTLY equal (eps, dims, shapes...). + - Pattern placeholders are wildcards (match any producer). + - Containment: every internal pattern node's users must be inside the match; an + intermediate consumed elsewhere in the graph (or returned by the model) cannot be + internal to your pattern. +- **A pattern must have exactly ONE returning value.** The harness builds the replacement + as a single opaque call node; `_replace_pattern` asserts + `len(match.returning_nodes) == len(copied_returning_nodes)` — a pattern returning + 2+ values CRASHES the run (AssertionError → variant fails). `return out` and + `return (out,)` are both fine (1 returning node). +- Replacement wiring: the harness wraps your pass as + `replacement(*pattern_args) = with_dispatch_wrapper_run(*replacement_args(*pattern_args))`. + `replacement_args` IS traced with FX proxies — it must be a pure argument shuffle + (reorder/drop/append constants like a route string; no tensor math, no `.shape` access). +- `replacement_func()` is called MULTIPLE times and must return the SAME module-level + function object each call (no lambdas/closures) — verified by `f() is f()`. +- **Global single replacement function**: there is ONE process-global + `g_replacement_func`; `set_g_replacement_func` asserts all passes share it, and + `output_pass_replacement_func_limit: 1` silently DROPS passes whose + `replacement_func()` differs from the (randomly) chosen one. Therefore with ≥2 pass + files, every pass must `from pass_dir._shared_kernels import dispatch_wrapper` and + return that same object; per-pass behaviour is selected by a constant route string + appended in `replacement_args`. (`sys.path` contains the sample root during pass + loading, so `pass_dir.xxx` imports work; no `__init__.py` needed.) +- Check the log line `[PassMgrBackend] Loaded N passes: [...]` — if N < your file count, + passes were dropped (validation failure or replacement-func limit). +- A pass listed in the JSON but failing AST validation is skipped with + "Detected hacking behavior ... Skipping loading". + +## 4. Anti-cheat (two layers — do not fight them, design within them) + +**Layer 1 — AST validation** (`pass_bench/ast_util/validate_pass_source.py`), applied to +every file listed in `sorted_output_pass_rule_names.json`: +- Blocked imports ANYWHERE in the file: `torch.nn`, `torch.nn.functional`, `torch.ops`, + `torch.autograd` (including `F = torch.nn.functional` alias assignments). Inside + `pattern()` use full dotted names: `torch.nn.functional.relu(...)`. +- In all code EXCEPT the bodies of `pattern` and `replacement_args`: any call resolving to + `torch.*` is blocked except `torch.empty, torch.empty_like, torch.zeros, + torch.zeros_like, torch.ones, torch.ones_like, torch.full, torch.full_like, + torch.as_tensor`. Tensor METHOD calls pass the AST check but are stopped by layer 2. +- A custom checker hook (`PASSNET_CUSTOM_PASS_SOURCE_CHECKER_PATH`) may add stricter rules + on the leaderboard. Keep helper modules (e.g. `_shared_kernels.py`) clean too, even + though the loader currently only validates listed pass files. + +**Layer 2 — Poison dispatch** (`posion_dispatch_tensor.py`): during the 25 warmup calls, +your wrapper receives `PosionDispatchTensor` args. Allowed on them: the factory ops above, +`.to()` (`aten._to_copy`), `torch.as_tensor`, and metadata (`.shape`, `.stride()`, +`.numel()`, `.dim()`, `.device`, `.dtype`, `.data_ptr()`). EVERY other aten op +(`.contiguous()`, `.view()`, `+`, `.relu()`, `torch.matmul`...) raises +"Unauthorized Operator" → the variant fails. Triton kernel launches work (they only read +`.data_ptr()` and metadata). Timed trials run unpoisoned, but a wrapper that only works +unpoisoned never survives warmup. + +Consequences: +- ALL math happens inside `@triton.jit` kernels. +- You cannot `.contiguous()` inputs — pass strides into the kernel instead. +- Allocate outputs with `torch.empty(...)`/`torch.empty_like(...)`. +- `x.to(dtype)` is legal if you must unify dtypes (it costs a copy). + +**Known-cheat patterns that MUST NOT be used** (they appear in some legacy passes; they are +detected as hacking and/or are slower than eager anyway): `getattr(torch, "conv2d")` +laundering, `torch.utils._mode_utils.no_dispatch()`, monkeypatching harness modules, +writing to harness state, returning cached eager results. + +## 5. Timing/overhead facts (measured, A100, torch 2.7.1, triton 3.3.1) + +| item | cost | +|---|---| +| Triton `kernel[grid](...)` Python launch | ~19 µs/call | +| aten elementwise op (e.g. relu) e2e | ~8 µs | +| aten batch_norm (eval mode) e2e | ~24 µs | +| `torch.empty_like` | ~2.5 µs | +| compiled-vs-eager fixed tax (guards+FX+wrapper+launch) | ~40–70 µs/call | + +Implication: replacing ONE cheap op loses (measured 0.77× on a 3-op graph); win by +absorbing MANY ops per launch or replacing genuinely expensive regions. On large tensors +(ms-scale kernels) the tax is negligible and kernel quality dominates. + +## 6. Input replay (`pass_bench/torch/utils.py::replay_tensor`) + +For each entry in `weight_meta.py`: if `data` present → exact values; else +`randn(shape)·std·0.2 + mean` (std=0 → constant tensor), clamped to `[min_val, max_val]` +if given, non-finite→small noise, clamped to [-100, 100], then `.to(dtype).to(device)`. +Inputs are CONTIGUOUS, freshly replayed; `device` in meta may say `cpu` but everything is +moved to cuda. Model is called with inputs bound by forward-signature parameter NAMES. +Use mean/std/min_val to reason about numerics (e.g. running_var ≥ 0, masks are 0/1). + +## 7. Evaluation modes (how to read/write/evaluate) + +**Local mode** (you are in a sample dir; no HTTP service): +```bash +ls graph_list.txt pass_dir/ && bash entry.sh # full eval; log + score under $PASSNET_EVAL_OUTPUT (default /tmp/workspace_pass_bench_test) +``` + +**Sample Access Service mode** (`/health` returns `"mode": "sample_access_service"`), +multi-tenant; `SVC` and `SAMPLE` are given in your task prompt; every request needs +`?sample_path=$SAMPLE`: +```bash +curl -s "$SVC/problem?sample_path=$SAMPLE" # graph_list + model_code + weight_meta per graph +curl -s "$SVC/files?sample_path=$SAMPLE" # list pass_dir files +curl -s "$SVC/files/X.py?sample_path=$SAMPLE" # read +python3 -c 'import json,sys; print(json.dumps(dict(content=sys.stdin.read())))' < local/X.py | \ + curl -s -X POST "$SVC/files/X.py?sample_path=$SAMPLE" -H 'Content-Type: application/json' -d @- +curl -s -X DELETE "$SVC/files/X.py?sample_path=$SAMPLE" +curl -s -X POST "$SVC/evaluate" -H 'Content-Type: application/json' \ + -d "{\"sample_path\":\"$SAMPLE\"}" --max-time 610 # {"returncode","pass_matched","score","stdout","stderr"} +``` +Only `pass_dir/*.py|*.json` are writable; 503 = GPU busy → wait 15 s, retry (≤5×); +eval timeout 600 s. The returned `stdout` is pre-filtered (Trial/allclose lines removed) — +save it to a file and run `parse_eval_log.py` on it. + +**API-server mode** (single-sample server): same endpoints WITHOUT `sample_path` +(`/problem`, `/files/...`, `POST /evaluate`). + +## 8. Misc facts that bite + +- `model.py` lines like `tmp_4 = None` are refcount hints, NOT graph nodes — never put + them in a pattern. +- A sample's `graph_list.txt` typically holds 2–10 variants but can hold 100+; eval time + and timeout risk scale with it. +- The eager model output dtype must equal yours exactly (`[Datatype]` check) — store in + the input/output dtype, never leave fp32 results when eager returns fp16. +- The same pass set runs against ALL variants; per-variant "pass A matches, pass B + doesn't" is fine as long as ≥1 pass matches per variant (each variant only needs the + graph to be modified by at least one pass). +- Subprocess-per-variant means module-level state in your pass file does NOT persist + across variants. +- `aggregated_score.json` is written per run; the service `score` field echoes it + (averaged over `num_runs` when the service evaluates multiple times). diff --git a/skills/task-oriented/Passnet/skills/passnet-triton-opt/SKILL.md b/skills/task-oriented/Passnet/skills/passnet-triton-opt/SKILL.md new file mode 100644 index 000000000..7753142d3 --- /dev/null +++ b/skills/task-oriented/Passnet/skills/passnet-triton-opt/SKILL.md @@ -0,0 +1,181 @@ +--- +name: passnet-triton-opt +description: > + Make a single PassNet Triton kernel fast AND numerically faithful: performance model + (when a replacement can win at all), block/grid/warp tuning, autotune policy, launch + overhead, and per-op numeric recipes to pass the dtype baseline tolerances. Use when a + pass matches and is correct but speedup below expected, or when correctness fails by small + numeric margins. +--- + +## 1. Performance model — know if a win is even possible + +Measured costs (A100, torch 2.7.1 / triton 3.3.1; same order of magnitude on the eval box): + +| component | cost per call | +|---|---| +| dynamo guards + FX graph interpretation + dispatch wrapper | ~25–45 µs on a 3-op graph; grows with node count (~150–230 µs seen on a 13-node graph) — calibrate from your first eval | +| Triton kernel Python launch (`kernel[grid](...)`) | ~19 µs each | +| `torch.empty_like` / `torch.empty` | ~2–3 µs each | +| each aten op you REMOVED from the graph | +5–25 µs back (small ops) or its real kernel time | + +So per variant: `compiled_e2e ≈ eager_e2e − Σ(absorbed aten µs) + tax + Σ(your kernels µs)`. +Rules of thumb: +- Whole eager forward < ~150 µs and you can absorb ≤2 cheap ops → ceiling < 1. Ship a + safe correct pass, accept ~0.8–0.95, spend your time on another bottleneck. +- Small or narrow fusions can match and be numerically correct yet still lose e2e when + launch/framework overhead exceeds the eager work removed. +- Layout-only, view/reshape-only, split/fanout, tiny activation, and small reduction-like + kernels can also lose even when their local kernel is correct: eager may treat the work + as metadata, use a vendor-tuned primitive, or pay less framework overhead than the + compiled path. Treat these as floor/downside caps unless they bridge substantial compute + or remove an actual materialization. +- The lever is **ops absorbed per launch** (and avoiding extra launches/allocs), not + shaving 10% off an already-tiny kernel. +- Big tensors (≥ a few MB; kernel ≥ 100 µs): overhead noise vanishes; now memory traffic + is everything — a clean fused kernel that reads each input once and writes once usually + beats the multi-pass eager chain by the ratio of memory passes saved. +- **The biggest wins can be ALGORITHMIC, not tuning** — but the default for a heavy conv/matmul + is still "leave it in the library, fuse its tail," because the library is near-optimal inside + the regime it tuned for. The clearest case where a rewrite beats it: a **non-overlapping + convolution** (`stride == kernel_size`, `kernel_size > 1`, `padding == 0`) is a dense patch + matmul while cuDNN runs a heavier general im2col path — a `tl.dot` patch-gather kernel + (kernel-templates §13) can win big. It usually does NOT pay to rewrite a **1×1 conv** (already + a cuBLAS-optimal pointwise GEMM), an **overlapping conv** (im2col blows up memory), a + **depthwise/grouped conv**, or a **general dense matmul** — those are vendor sweet spots. + Treat this as a prior, not a ban: the completed eval decides. Pursue any rewrite only after a + floor is banked, and don't dismiss the non-overlapping case as "cuDNN, unbeatable" without + checking `stride == kernel_size`. +- For large cat/stack/gather-style outputs, one giant kernel is not automatically best if it + needs heavy branching or irregular index logic. A small fixed number of simple kernels can + be faster and more timing-stable when each launch writes a large contiguous region. Compare + sequential micro-bench signals, then trust the completed evaluation. +- Repeated round evidence favors long memory-bound elementwise and broadcast-affine chains + when one legal single-output kernel absorbs multiple affine, activation, residual, layout, + or normalization consumers. These are the first places to spend tuning effort after a + correct floor because saved memory passes can amortize wrapper and launch overhead. +- Repeated completed evaluations also favor normalization-affine compute prefixes and + concat-to-normalization-affine prefixes when they remove large memory passes in one launch. + Treat kernel layout choices for concat branches as workload-dependent: a broad prefix can + win strongly, but small/branchy cases or attempts to absorb extra producers can regress in + score, correctness, or timing stability. Keep the best completed floor/prefix state and + revert completed regressions instead of stacking more small kernels or sweeping blindly. +- Reduction-only regions and vendor-tail replacements still need completed-evaluation proof. + A correct reduction kernel, a small tail behind a vendor op, or a region that replaces a + vendor-tuned primitive can remain below eager even when local kernel timing looks plausible. + Use the completed eval to decide keep/revert; do not infer success from the triage family. +- Compare `[Speedup][gpu]` vs `[Speedup][e2e]` in the eval log: gpu≈e2e&both<1 → on-stream + gaps/launches dominate (absorb more ops / fewer launches); gpu>1 but e2e<1 → host-side + overhead (reduce launches, drop autotune key churn, simplify wrapper Python). +- Only a completed evaluation can prove the final overhead-bound conclusion. Local + micro-bench and `check_pattern --bench` are pre-flight signals; if the completed eval is + matched and correct but below eager for the maximal legal region, stop or revert instead + of sweeping launch-bound tuning knobs. + +## 2. Kernel quality checklist (in order of impact) + +1. **One read per input element, one write per output element.** No multi-launch chains + when one kernel can hold the chain in registers. Exception: for heterogeneous + direct-output cat/stack branches, compare a branch-heavy single kernel with grouped + launches; choose by measured stable e2e speed, not by launch count alone. +2. **Contiguous innermost access**: thread `offs` should walk the last (stride-1) dim. + For strided/permuted inputs pass strides; never `.contiguous()` (poison-illegal) — + absorbing the transpose INTO the index math is free. +3. **Block sizing**: elementwise → BLOCK 1024 (n ≥ 1M: 2048, num_warps=8). Row kernels → + `BLOCK_N = next_power_of_2(N)`, and start with **num_warps=4** — more warps on small + rows HURTS (measured: 512-elem rows, warps=8 was 1.8× slower than warps=4). If + N > 16384 switch to a looped accumulation (template 6 in kernel-templates). + For tiny rows (≤1024 elems), SWEEP rows-per-program ∈ {1, 2, 4, 8} (2D tile) — the + winner is shape-dependent (measured 1.45× for ROWS=4 at D=512/131k rows, but ROWS=1 + best on a 768-wide LayerNorm); never assume, always micro-bench (§5). +4. **Grid**: prefer 1D `cdiv(n, BLOCK)`; for row kernels one program per row is fine up + to ~256k rows. Avoid grids of 1–4 programs on big inputs (underutilization) — split. +5. **num_stages**: only matters for `tl.dot`/looped kernels (3–4); elementwise ignore. +6. **Avoid recompilation churn**: every distinct `tl.constexpr` value (incl. BLOCK chosen + from shape) compiles a new binary during warmup. 2–3 specializations fine; per-shape + formulas over 100-variant samples → timeout risk. Derive constexprs from a small fixed + set (e.g. BLOCK_N = next_pow2 clamped to {128, 512, 1024, 4096}). +7. **Autotune policy**: `@triton.autotune` is legal and runs during the untimed warmup + (poison-safe). Use ≤4–6 configs, `key=['N']`-style coarse keys. SKIP autotune when the + sample has many graph variants (tuning cost × variants × 600 s timeout) or when the + kernel is launch-bound anyway. A hand-picked config from one local micro-benchmark is + usually within 5%. +8. **Scalar args beat tensor args**: precompute `n`, `C`, `HW`, strides in the wrapper + (metadata access is poison-legal) — never load shape info from tensors in-kernel. +9. **Reduce wrapper Python**: no dict lookups in hot path beyond the route `if`; compute + grid inline; allocate with `torch.empty` (not zeros — don't pay a memset you'll + overwrite). + +## 3. Numeric fidelity — passing t=-5 baseline tolerances + +Baselines: fp32 rtol 1.3e-6/atol 1e-5 (tight — mirror eager's algorithm), fp16 rtol 1e-3, +bf16 rtol 1.6e-2 (loose). The same kernel must satisfy ALL dtype variants. + +Recipes (these match aten's CUDA implementations): +- **Always compute in fp32** (`.to(tl.float32)` after load), store with + `.to(out_ptr.dtype.element_ty)`. For fp16/bf16 inputs this reproduces aten (which also + accumulates in fp32) to ≤1 ulp. +- Reductions: `tl.sum`/`tl.max` over a row in one block ≈ aten's order closely enough for + fp32 rtol 1.3e-6 at typical sizes (N ≤ 16k). For looped accumulation keep ONE fp32 + accumulator (don't reorder into multiple partial schemes unnecessarily). +- softmax: subtract row max (aten does), exp, divide by sum. Don't use exp2 shortcuts on + fp32 variants. +- layer_norm/batch_norm: biased variance (÷N), `1/sqrt(var+eps)` in fp32, then scale/shift. +- mean: `sum/N` (single division at the end, not running mean). +- GELU exact: `0.5x(1+erf(x/√2))` via `tl.math.erf`; tanh GELU only when model.py says + `approximate='tanh'`. +- sigmoid/silu: `1/(1+exp(-x))` fp32 — fine at all baselines. +- `x.pow(2)` → `x*x`; `torch.rsqrt` → `1.0/tl.sqrt` (matches aten rsqrt within 1 ulp; if + a strict fp32 variant complains, try `tl.math.rsqrt`). +- Integer/bool tensors: compare with `torch.equal` (atol/rtol=0 conceptually) — your + kernel must be EXACT (indices, masks, argmax-style ops: replicate tie-breaking by + first-occurrence). +- Output dtype must equal eager's per output ([Datatype] check), including intermediate + `.to(torch.float32)` nodes you absorbed — store in the dtype eager would have at the + region's output. + +Diagnosing accuracy failures from eval logs: `[Correctness][max_diff]` magnitude vs the +table above (parse_eval_log.py does this) — +- max_diff ~1e-3·|values| on fp16 only → fp16 rounding mismatch: ensure fp32 internal math. +- fails ONLY fp32 variants → algorithmic mismatch (reduction order, formula) — mirror + aten's exact sequence. +- max_diff huge (≥1) → indexing/masking bug, not numerics: check strides, `other=` values + contaminating reductions (use `-inf` for max, `0` for sum, and mask BEFORE divide), + uninitialized output regions (mask coverage), `offs` overflow on >2^31 elements + (use `tl.int64` offsets when numel > 2e9). +- NaN/Inf only in compiled → division by masked-out zero lanes; `tl.where` the mask before + the division, not after. + +## 4. Launch-overhead reduction (when e2e-bound) + +- Merge kernels (one launch per region; the whole point of fusion). +- Do not add a tiny independent kernel after a correct-but-slow state unless it absorbs + enough real eager work to pay for another launch. If a completed eval shows the added + region lowers score, revert to the best completed state and stop that line. +- Avoid per-call `next_power_of_2` recompute churn → precompute in wrapper; constexpr set + small (see 2.6). +- Don't return tuples/lists from the wrapper unless the pattern output needs it. +- Buffer reuse across calls (cache `out` keyed by shape in a module dict) saves ~2 µs but + risks aliasing bugs if the model returns your output AND mutates downstream — only as a + last resort, verify with full eval. +- CUDA graphs / `torch.compile` of the wrapper / `no_dispatch` tricks: NOT allowed + (torch.* calls blocked / hacking). Don't. + +## 5. Micro-benchmark before full eval (GPU, seconds) + +```python +import torch, time +def bench(fn, *args, iters=500): + for _ in range(20): fn(*args) + torch.cuda.synchronize(); t0 = time.perf_counter() + for _ in range(iters): fn(*args) + torch.cuda.synchronize(); return (time.perf_counter() - t0) / iters * 1e6 # µs +``` +Compare: (a) eager region ops chained, (b) your wrapper. Build inputs with the shapes / +dtypes from weight_meta (use the checker's replay). Target: wrapper ≤ eager_region − 40 µs +for small graphs; for ms-scale kernels target the memory-traffic bound (bytes moved / +~1.5 TB/s on A100 ≈ achievable µs). + +`check_pattern.py --bench` (passnet-feedback) automates exactly this comparison per +variant; trust the full eval for the final number (it adds guards/FX costs that local +micro-bench misses).