From 3895b9d1829dd29d4b0c440b6dea1254c17ee91d Mon Sep 17 00:00:00 2001 From: zihaoan2 Date: Mon, 29 Jun 2026 05:01:52 +0000 Subject: [PATCH 1/5] feat(e2e/apply): native (.cu/.hip/.cpp) apply-back plumbing in overlay_setup.py [Phase 1] MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The e2e overlay is Python-only (sitecustomize/PYTHONPATH inject + attr rebind), so a compiled-source kernel rewrite (aiter/CK .cu, local HIP .hip, cpp_extension, ...) can be measured isolated but never deployed into the live server — it gets rejected no_rebind_seam/editable=false. This adds the framework-agnostic, reversible plumbing to apply such a patch IN PLACE and undo it byte-exact, while keeping ONE tool + ONE manifest + ONE revert shared with the existing Python overlay. This is the safety spine only; it never invents a build command — the caller (e2e_integrator agent, Phase 2) discovers the install's own incremental build the same way benchmark_engineer does and passes it via --build-cmd, so the design stays general and hardcode-free. overlay_setup.py — new manifest section "natives" + subcommands: - add-native: back up source + named artifacts + tar-snapshot the invalidated cache dir(s), save the manifest BEFORE mutating (crash-safe: every entry references an existing backup), then swap the source in place, delete the cache(s) to force an incremental rebuild, and run the caller-supplied build. Only the changed translation unit rebuilds (ninja/jit/cpp_itfs cache keyed by the change). - verify-native: confirm each tracked artifact actually changed/was created (catch a silent no-op build). - revert: unified — Python overlay drops off PYTHONPATH; native sources/artifacts/cache dirs restored byte-exact from backup, reverse order, idempotent (natives cleared only on full success). - gc-stale: scan a root for a crashed run's still-applied native overlay and revert it, so a fresh session never starts on a dirty install. test_overlay_native.py — stdlib unittest (no GPU/server): apply+rebuild, fresh-build verify (pass + silent-no-op fail), byte-exact revert (source/artifact/cache), artifact-created-then-removed, double-revert idempotency, revert-without-build, mixed Python-overlay + native round-trip, gc-stale. 9 tests green. Default manifest gains "natives": [] (backward-compatible; existing flows untouched). Co-Authored-By: Claude Opus 4.8 --- e2e_workflow/scripts/overlay_setup.py | 225 +++++++++++++++++++- e2e_workflow/scripts/test_overlay_native.py | 199 +++++++++++++++++ 2 files changed, 422 insertions(+), 2 deletions(-) create mode 100644 e2e_workflow/scripts/test_overlay_native.py diff --git a/e2e_workflow/scripts/overlay_setup.py b/e2e_workflow/scripts/overlay_setup.py index e0df37dc8..f09a10994 100755 --- a/e2e_workflow/scripts/overlay_setup.py +++ b/e2e_workflow/scripts/overlay_setup.py @@ -29,9 +29,28 @@ --module sglang.srt.layers.activation Back-compat aliases: `monkeypatch` == add-rebind, `copy-subtree` == add-module (file granularity). + +NATIVE apply-back (compiled-source kernels: .cu/.hip/.cpp/CK/...) — these CANNOT be deployed via the +non-invasive PYTHONPATH overlay above (they need a recompiled .so in the install). They are applied +IN PLACE to the install and tracked in the same manifest's "natives" section so a single `revert` +undoes both the PYTHONPATH overlay AND the in-place native changes. This helper only does the +safety-critical, framework-agnostic plumbing (back up → byte-exact restore → cache-dir snapshot → +fresh-build verify); it NEVER invents a build command — the caller (e2e_integrator agent) discovers +the install's own incremental build the same way benchmark_engineer does and passes it via --build-cmd. + + add-native apply a compiled-source patch in place + (optionally) run an incremental rebuild + --overlay O --target (--patched-file F | --patch D) + [--artifact A ...] [--invalidate-cache DIR ...] [--build-cmd "..."] [--build-cwd C] + verify-native confirm each tracked artifact actually changed (fresh build happened, not a silent no-op) + --overlay O + revert restore EVERYTHING (PYTHONPATH overlay is just dropped; native files/artifacts/caches + restored byte-exact from backup). Idempotent. --overlay O + gc-stale scan a root for native overlays left applied by a crashed run and revert them + --root R + Stdlib only. """ -import argparse, importlib, json, os, shutil, subprocess, sys +import argparse, glob, hashlib, importlib, json, os, shlex, shutil, subprocess, sys, tarfile SITECUSTOMIZE = r'''# Auto-generated reversible overlay (e2e_workflow). Drop this dir from PYTHONPATH to revert. import json, os, sys, importlib, importlib.util @@ -112,7 +131,7 @@ def _ensure_overlay(overlay): man = os.path.join(overlay, "_overlay_manifest.json") if not os.path.exists(man): with open(man, "w") as fh: - json.dump({"modules": [], "rebinds": [], "captures": []}, fh, indent=2) + json.dump({"modules": [], "rebinds": [], "captures": [], "natives": []}, fh, indent=2) return man @@ -198,6 +217,184 @@ def cmd_check(a): ("INJECTED" if f.endswith(a.module + ".py") else "INSTALL (overlay not shadowing this module)")) +# ---------------------------------------------------------------------------------------------------- +# NATIVE apply-back: in-place, reversible deploy of a COMPILED-source kernel patch. +# Framework-agnostic plumbing only — the caller discovers + passes the incremental build command. +# ---------------------------------------------------------------------------------------------------- +def _sha256(path): + h = hashlib.sha256() + with open(path, "rb") as fh: + for chunk in iter(lambda: fh.read(1 << 20), b""): + h.update(chunk) + return h.hexdigest() + + +def _native_backup_dir(overlay): + d = os.path.join(overlay, "_native_backup") + os.makedirs(d, exist_ok=True) + return d + + +def cmd_add_native(a): + """Apply a compiled-source patch in place (reversibly), optionally invalidating caches + rebuilding. + + Crash-safe ordering: every backup is created AND the manifest entry saved BEFORE the install is + mutated, so a crash at any point leaves a manifest whose backups all exist -> revert/gc-stale can + always restore byte-exact. Per-item restore is idempotent. + """ + man = _ensure_overlay(a.overlay) + bdir = _native_backup_dir(a.overlay) + m = _load_man(man) + m.setdefault("natives", []) + n = len(m["natives"]) + + tgt = os.path.abspath(a.target) + if not os.path.exists(tgt): + raise SystemExit(f"add-native: target source does not exist: {tgt}") + if not (a.patched_file or a.patch): + raise SystemExit("add-native: requires --patched-file or --patch") + + entry = {"sources": [], "artifacts": [], "caches": [], + "build_cmds": [], "build_cwd": a.build_cwd or "", "verify": {}} + + # 1. back up the source FIRST (so the manifest never references a missing backup) + src_bak = os.path.join(bdir, f"src_{n}_{os.path.basename(tgt)}") + shutil.copy2(tgt, src_bak) + entry["sources"].append({"install_path": tgt, + "backup": os.path.relpath(src_bak, a.overlay), + "sha256": _sha256(tgt)}) + + # 2. back up named artifacts (pre-build); the first existing one is the fresh-build verify anchor + for i, art in enumerate(a.artifact or []): + art = os.path.abspath(art) + if os.path.exists(art): + art_bak = os.path.join(bdir, f"art_{n}_{i}_{os.path.basename(art)}") + shutil.copy2(art, art_bak) + sha = _sha256(art) + entry["artifacts"].append({"path": art, "backup": os.path.relpath(art_bak, a.overlay), + "sha256": sha, "existed": True}) + if not entry["verify"]: + entry["verify"] = {"artifact": art, "pre_sha256": sha} + else: + entry["artifacts"].append({"path": art, "backup": "", "sha256": "", "existed": False}) + + # 3. snapshot (tar) each cache dir we will invalidate, so revert restores it intact + for i, cd in enumerate(a.invalidate_cache or []): + cd = os.path.abspath(cd) + existed = os.path.isdir(cd) + rel_tar = "" + if existed: + tar_bak = os.path.join(bdir, f"cache_{n}_{i}.tar") + with tarfile.open(tar_bak, "w") as tf: + tf.add(cd, arcname=os.path.basename(cd)) + rel_tar = os.path.relpath(tar_bak, a.overlay) + entry["caches"].append({"dir": cd, "backup_tar": rel_tar, "existed": existed}) + + if a.build_cmd: + entry["build_cmds"].append(shlex.split(a.build_cmd)) + + # SAVE manifest BEFORE any mutation (the crash-safety point) + m["natives"].append(entry) + _save_man(man, m) + + # 4. MUTATE the install: write the patched source in place + if a.patched_file: + shutil.copy2(a.patched_file, tgt) + elif a.patch and not _try_apply(a.patch, target_file=tgt): + shutil.copy2(src_bak, tgt) # roll back the one file we touched, then abort + raise SystemExit(f"add-native: failed to apply patch {a.patch} to {tgt}") + + # delete invalidated caches to force a rebuild + for c in entry["caches"]: + if c["existed"] and os.path.isdir(c["dir"]): + shutil.rmtree(c["dir"]) + + # 5. run the caller-supplied incremental build (we never invent it) + for cmd in entry["build_cmds"]: + r = subprocess.run(cmd, cwd=(a.build_cwd or None), capture_output=True, text=True) + sys.stderr.write(r.stdout + r.stderr) + if r.returncode != 0: + raise SystemExit(f"add-native: build failed ({' '.join(cmd)}) rc={r.returncode}") + + print(f"OVERLAY_DIR={a.overlay}") + print(f"add-native {tgt} (sources=1 artifacts={len(entry['artifacts'])} caches={len(entry['caches'])} build={'yes' if entry['build_cmds'] else 'no'})") + + +def cmd_verify_native(a): + """Confirm each tracked artifact actually changed (or was created) — catches a silent no-op build.""" + man = os.path.join(a.overlay, "_overlay_manifest.json") + if not os.path.exists(man): + raise SystemExit("verify-native: no manifest") + m = _load_man(man) + ok, checked = True, 0 + for entry in m.get("natives", []): + for art in entry.get("artifacts", []): + p = art["path"] + if not os.path.exists(p): + print(f"FRESH_BUILD_FAIL {'missing-after-build' if art.get('existed') else 'not-created'} {p}") + ok = False + continue + if art.get("existed") and _sha256(p) == art.get("sha256"): + print(f"FRESH_BUILD_FAIL unchanged {p}") + ok = False + else: + checked += 1 + print(("FRESH_BUILD_OK" if ok else "FRESH_BUILD_FAIL") + f" (checked {checked})") + if not ok: + raise SystemExit(2) + + +def cmd_revert(a): + """Restore everything. PYTHONPATH overlay reverts by simply not being on PYTHONPATH; native changes + are restored byte-exact from backup (sources + artifacts + cache dirs), in reverse apply order. + Idempotent: each item restores from its own backup, and the natives list is cleared only on full success.""" + man = os.path.join(a.overlay, "_overlay_manifest.json") + if not os.path.exists(man): + print("revert: no manifest; nothing to do") + return + m = _load_man(man) + for entry in reversed(m.get("natives", [])): + for c in entry.get("caches", []): + d = c["dir"] + if os.path.isdir(d): + shutil.rmtree(d) + if c.get("existed") and c.get("backup_tar"): + tarp = os.path.join(a.overlay, c["backup_tar"]) + if os.path.exists(tarp): + with tarfile.open(tarp) as tf: + tf.extractall(os.path.dirname(d)) + # existed=False -> candidate-built cache: leave it removed + for art in entry.get("artifacts", []): + if art.get("existed") and art.get("backup"): + bak = os.path.join(a.overlay, art["backup"]) + if os.path.exists(bak): + shutil.copy2(bak, art["path"]) + elif not art.get("existed") and os.path.exists(art["path"]): + os.remove(art["path"]) # candidate created it -> remove + for s in entry.get("sources", []): + bak = os.path.join(a.overlay, s["backup"]) + if os.path.exists(bak): + shutil.copy2(bak, s["install_path"]) + m["natives"] = [] + _save_man(man, m) + print(f"revert: restored native changes in {a.overlay}") + + +def cmd_gc_stale(a): + """Scan a root for native overlays left applied by a crashed run (manifest has a non-empty natives + list) and revert each — so a fresh session never starts on a dirty install.""" + found = 0 + for man in glob.glob(os.path.join(a.root, "**", "_overlay_manifest.json"), recursive=True): + try: + m = _load_man(man) + except Exception: + continue + if m.get("natives"): + cmd_revert(argparse.Namespace(overlay=os.path.dirname(man))) + found += 1 + print(f"gc-stale: reverted {found} stale native overlay(s) under {a.root}") + + def main(): ap = argparse.ArgumentParser() sub = ap.add_subparsers(dest="cmd", required=True) @@ -235,6 +432,30 @@ def main(): p.add_argument("--module", required=True) p.set_defaults(func=cmd_check) + p = sub.add_parser("add-native") + p.add_argument("--overlay", required=True) + p.add_argument("--target", required=True, help="absolute path of the install source file to patch") + p.add_argument("--patched-file", default="", dest="patched_file", help="full replacement source") + p.add_argument("--patch", default="", help="unified diff to apply to --target") + p.add_argument("--artifact", action="append", default=[], help="built artifact(s) to snapshot for revert/verify (repeatable)") + p.add_argument("--invalidate-cache", action="append", default=[], dest="invalidate_cache", + help="cache dir(s) to snapshot+delete so only the changed unit rebuilds (repeatable)") + p.add_argument("--build-cmd", default="", dest="build_cmd", help="caller-discovered incremental build command") + p.add_argument("--build-cwd", default="", dest="build_cwd") + p.set_defaults(func=cmd_add_native) + + p = sub.add_parser("verify-native") + p.add_argument("--overlay", required=True) + p.set_defaults(func=cmd_verify_native) + + p = sub.add_parser("revert") + p.add_argument("--overlay", required=True) + p.set_defaults(func=cmd_revert) + + p = sub.add_parser("gc-stale") + p.add_argument("--root", required=True) + p.set_defaults(func=cmd_gc_stale) + a = ap.parse_args() a.func(a) diff --git a/e2e_workflow/scripts/test_overlay_native.py b/e2e_workflow/scripts/test_overlay_native.py new file mode 100644 index 000000000..95af0df45 --- /dev/null +++ b/e2e_workflow/scripts/test_overlay_native.py @@ -0,0 +1,199 @@ +#!/usr/bin/env python3 +"""Stdlib unittest for overlay_setup.py's NATIVE apply-back plumbing (no GPU, no server, CI-safe). + +Covers the safety-critical guarantees: in-place apply + byte-exact revert (sources, named artifacts, +cache dirs), fresh-build verification, mixed Python-overlay + native round-trip, crash-recovery +idempotency, and gc-stale cleanup of a crashed run's dirty install. + +Run: python3 test_overlay_native.py +""" +import hashlib +import json +import os +import shutil +import subprocess +import sys +import tempfile +import unittest + +HERE = os.path.dirname(os.path.abspath(__file__)) +SCRIPT = os.path.join(HERE, "overlay_setup.py") + + +def sha(path): + h = hashlib.sha256() + with open(path, "rb") as fh: + h.update(fh.read()) + return h.hexdigest() + + +def rt(path): + with open(path) as fh: + return fh.read() + + +def rb(path): + with open(path, "rb") as fh: + return fh.read() + + +def run(*args, check=True): + r = subprocess.run([sys.executable, SCRIPT, *args], capture_output=True, text=True) + if check and r.returncode != 0: + raise AssertionError(f"cmd failed ({r.returncode}): {' '.join(args)}\n{r.stdout}\n{r.stderr}") + return r + + +class NativeApplyBack(unittest.TestCase): + def setUp(self): + self.tmp = tempfile.mkdtemp(prefix="ovltest_") + # a fake "install" tree: a compiled source + its prebuilt artifact + a regenerable cache dir + self.install = os.path.join(self.tmp, "install", "pkg", "kernels") + os.makedirs(self.install) + self.src = os.path.join(self.install, "gemm.cu") + self.art = os.path.join(self.install, "gemm.so") + self.cache = os.path.join(self.install, ".jit_cache", "gemm_abc123") + os.makedirs(self.cache) + with open(self.src, "w") as fh: + fh.write("// ORIGINAL kernel source\n__global__ void gemm() {}\n") + with open(self.art, "wb") as fh: + fh.write(b"PREBUILT-ARTIFACT-ORIGINAL") + with open(os.path.join(self.cache, "compiled.hsaco"), "wb") as fh: + fh.write(b"OLD-CACHE-ENTRY") + self.src_orig_sha = sha(self.src) + self.art_orig_sha = sha(self.art) + + # the patched source the optimizer produced + self.patched = os.path.join(self.tmp, "gemm_patched.cu") + with open(self.patched, "w") as fh: + fh.write("// OPTIMIZED kernel source\n__global__ void gemm() { /*fast*/ }\n") + + # a stub "incremental build": rewrites the artifact (simulates a real rebuild touching the .so) + self.builder = os.path.join(self.tmp, "build.sh") + with open(self.builder, "w") as fh: + fh.write("#!/bin/bash\nprintf 'REBUILT-FROM-%s' \"$(cat '%s')\" > '%s'\n" + % ("OPTIMIZED", self.src, self.art)) + os.chmod(self.builder, 0o755) + + self.overlay = os.path.join(self.tmp, "overlay") + + def tearDown(self): + shutil.rmtree(self.tmp, ignore_errors=True) + + def manifest(self): + with open(os.path.join(self.overlay, "_overlay_manifest.json")) as fh: + return json.load(fh) + + # ---- core apply + build ---------------------------------------------------------------------- + def test_apply_swaps_source_and_rebuilds(self): + run("add-native", "--overlay", self.overlay, "--target", self.src, + "--patched-file", self.patched, "--artifact", self.art, + "--invalidate-cache", self.cache, + "--build-cmd", f"bash {self.builder}") + # source swapped in place + self.assertIn("OPTIMIZED", rt(self.src)) + # artifact rebuilt (changed) + self.assertNotEqual(sha(self.art), self.art_orig_sha) + self.assertIn("REBUILT", rb(self.art).decode()) + # manifest records one native entry with backups + nat = self.manifest()["natives"] + self.assertEqual(len(nat), 1) + self.assertEqual(nat[0]["sources"][0]["sha256"], self.src_orig_sha) + self.assertTrue(nat[0]["verify"]["artifact"].endswith("gemm.so")) + + # ---- fresh-build verification ---------------------------------------------------------------- + def test_verify_passes_when_artifact_changed(self): + run("add-native", "--overlay", self.overlay, "--target", self.src, + "--patched-file", self.patched, "--artifact", self.art, + "--build-cmd", f"bash {self.builder}") + r = run("verify-native", "--overlay", self.overlay) + self.assertIn("FRESH_BUILD_OK", r.stdout) + + def test_verify_fails_on_silent_noop_build(self): + # a build that does NOT touch the artifact -> verify must catch it + noop = os.path.join(self.tmp, "noop.sh") + with open(noop, "w") as fh: + fh.write("#!/bin/bash\ntrue\n") + os.chmod(noop, 0o755) + run("add-native", "--overlay", self.overlay, "--target", self.src, + "--patched-file", self.patched, "--artifact", self.art, + "--build-cmd", f"bash {noop}") + r = run("verify-native", "--overlay", self.overlay, check=False) + self.assertNotEqual(r.returncode, 0) + self.assertIn("FRESH_BUILD_FAIL", r.stdout) + + # ---- byte-exact revert ----------------------------------------------------------------------- + def test_revert_restores_everything_byte_exact(self): + run("add-native", "--overlay", self.overlay, "--target", self.src, + "--patched-file", self.patched, "--artifact", self.art, + "--invalidate-cache", self.cache, "--build-cmd", f"bash {self.builder}") + run("revert", "--overlay", self.overlay) + # source restored byte-exact + self.assertEqual(sha(self.src), self.src_orig_sha) + # artifact restored byte-exact + self.assertEqual(sha(self.art), self.art_orig_sha) + # cache dir restored intact + self.assertTrue(os.path.isfile(os.path.join(self.cache, "compiled.hsaco"))) + self.assertEqual(rb(os.path.join(self.cache, "compiled.hsaco")), b"OLD-CACHE-ENTRY") + # natives list cleared + self.assertEqual(self.manifest()["natives"], []) + + def test_revert_removes_artifact_that_did_not_exist(self): + os.remove(self.art) # candidate build will CREATE it; revert must remove it + run("add-native", "--overlay", self.overlay, "--target", self.src, + "--patched-file", self.patched, "--artifact", self.art, + "--build-cmd", f"bash {self.builder}") + self.assertTrue(os.path.exists(self.art)) + run("revert", "--overlay", self.overlay) + self.assertFalse(os.path.exists(self.art)) + self.assertEqual(sha(self.src), self.src_orig_sha) + + # ---- crash-recovery idempotency -------------------------------------------------------------- + def test_double_revert_is_noop(self): + run("add-native", "--overlay", self.overlay, "--target", self.src, + "--patched-file", self.patched, "--artifact", self.art, "--build-cmd", f"bash {self.builder}") + run("revert", "--overlay", self.overlay) + run("revert", "--overlay", self.overlay) # must not raise / must keep things restored + self.assertEqual(sha(self.src), self.src_orig_sha) + self.assertEqual(sha(self.art), self.art_orig_sha) + + def test_revert_works_even_if_build_never_ran(self): + # simulate a crash right after manifest+mutate but before any verify: apply with no build cmd + run("add-native", "--overlay", self.overlay, "--target", self.src, + "--patched-file", self.patched, "--artifact", self.art, "--invalidate-cache", self.cache) + self.assertIn("OPTIMIZED", rt(self.src)) # mutated + run("revert", "--overlay", self.overlay) + self.assertEqual(sha(self.src), self.src_orig_sha) + self.assertTrue(os.path.isfile(os.path.join(self.cache, "compiled.hsaco"))) + + # ---- mixed Python-overlay + native ----------------------------------------------------------- + def test_mixed_python_and_native_same_overlay(self): + # a python rebind (non-invasive) + a native apply (invasive) share ONE overlay/manifest + run("add-rebind", "--overlay", self.overlay, + "--target", "pkg.mod:fn", "--impl-module", "impl", "--impl-attr", "fast_fn") + run("add-native", "--overlay", self.overlay, "--target", self.src, + "--patched-file", self.patched, "--artifact", self.art, "--build-cmd", f"bash {self.builder}") + man = self.manifest() + self.assertEqual(len(man["rebinds"]), 1) + self.assertEqual(len(man["natives"]), 1) + # one unified revert undoes the native (invasive) part; python overlay just stops being on PYTHONPATH + run("revert", "--overlay", self.overlay) + self.assertEqual(sha(self.src), self.src_orig_sha) + self.assertEqual(sha(self.art), self.art_orig_sha) + self.assertEqual(self.manifest()["natives"], []) + self.assertEqual(len(self.manifest()["rebinds"]), 1) # python overlay entry untouched + + # ---- gc-stale (crashed run left a dirty install) --------------------------------------------- + def test_gc_stale_reverts_crashed_run(self): + run("add-native", "--overlay", self.overlay, "--target", self.src, + "--patched-file", self.patched, "--artifact", self.art, "--build-cmd", f"bash {self.builder}") + # the run "crashed" — never called revert; install is dirty + self.assertIn("OPTIMIZED", rt(self.src)) + r = run("gc-stale", "--root", self.tmp) + self.assertIn("reverted 1", r.stdout) + self.assertEqual(sha(self.src), self.src_orig_sha) + self.assertEqual(sha(self.art), self.art_orig_sha) + + +if __name__ == "__main__": + unittest.main(verbosity=2) From e0efdc5588ba8d67fb2d35010720b68aeb7cdf01 Mon Sep 17 00:00:00 2001 From: zihaoan2 Date: Tue, 30 Jun 2026 04:04:56 +0000 Subject: [PATCH 2/5] feat(e2e): wire native (.cu/.hip/.cpp) apply-back into the e2e pipeline [Phase 2] MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Phase 1 added the framework-agnostic reversible plumbing (overlay_setup.py add-native/verify-native/ revert/gc-stale). Phase 2 wires the "brain" so a compiled-kernel win actually deploys end-to-end: - roles/e2e_integrator.md: new NATIVE patch path in the winner_kind=patch case. Detect native from apply_kind=="native" OR a code_patch touching a compiled suffix (.cu/.cuh/.hip/.cpp/.cc/.cxx/.c/.h/.hpp). DISCOVER the install's own incremental build (config.yaml compile_command / task_runner / Makefile / build.sh / ninja / CMake / torch cpp_extension load(name=) / aiter cpp_itfs MD_NAME) by walking up from the source — never invent it — and hand it to add-native via --build-cmd/--invalidate-cache. verify-native to catch silent no-op builds. A/B order is REF-first (clean install) -> add-native -> CAND, since native mutates the install globally; ALWAYS revert in a finally (accept/reject/crash) + assert install clean; serialize. Graph-safe + mem-footprint gates apply as for authored kernels. no_build_seam when no rebuildable seam exists. - roles/kernel_extractor.md: native (compiled-source) kernels in a rebuildable tree are now editable=true with apply_kind:"native" (target_callable:""); editable=false reserved for opaque prebuilt libs / read-only installs. .py/Triton stay apply_kind:"python" (unchanged). - e2e_workflow.js: add apply_kind ('python'|'native', default python) to EXTRACT_OP_SCHEMA + EXTRACT_SCHEMA; thread it into the serial-head and milestone KERNEL_RESULT inputs (the integrator also auto-detects native from patch suffixes, so missing threading degrades safely). - knowledge/learned/native-apply-back.md (+ INDEX): the method card. Default behavior byte-identical for Python/Triton winners (apply_kind defaults to python; the native path only triggers on a compiled-suffix patch or explicit apply_kind=native). node --check clean; Phase-1 unit tests still 9/9. Co-Authored-By: Claude Opus 4.8 --- e2e_workflow/e2e_workflow.js | 4 + e2e_workflow/knowledge/learned/INDEX.md | 1 + .../knowledge/learned/native-apply-back.md | 42 +++++++++++ e2e_workflow/roles/e2e_integrator.md | 73 +++++++++++++++++-- e2e_workflow/roles/kernel_extractor.md | 9 +++ 5 files changed, 122 insertions(+), 7 deletions(-) create mode 100644 e2e_workflow/knowledge/learned/native-apply-back.md diff --git a/e2e_workflow/e2e_workflow.js b/e2e_workflow/e2e_workflow.js index 265fb203f..05560d6f7 100644 --- a/e2e_workflow/e2e_workflow.js +++ b/e2e_workflow/e2e_workflow.js @@ -367,6 +367,7 @@ const EXTRACT_OP_SCHEMA = obj({ dtype: { type: 'string' }, synthesized: { type: 'boolean' }, regimes_captured: arrStr, candidate_backends: arrStr, reference_io_sha256: { type: 'string' }, target_callable: { type: 'string' }, // module:attr rebind seam for an authored kernel ('' if none) + apply_kind: { type: 'string' }, // 'python' (default) | 'native' (compiled-source kernel -> add-native in-place recompile) smoke: { type: 'string' }, notes: { type: 'string' }, }, ['op_kind', 'task_dir', 'smoke']); @@ -387,6 +388,7 @@ const EXTRACT_SCHEMA = obj({ source_path_in_sglang: { type: 'string' }, target_callable: { type: 'string' }, num_cases: { type: 'number' }, regimes_captured: arrStr, candidate_backends: arrStr, build: { type: 'boolean' }, unittest_smoke: { type: 'string' }, + apply_kind: { type: 'string' }, // 'python' (default) | 'native' reference_io_sha256: { type: 'string' }, notes: { type: 'string' }, }, ['editable', 'task_dir', 'unittest_smoke']); @@ -1601,6 +1603,7 @@ if (want('head') && headQueue.length && HEAD_BUDGET > 0) { KERNEL_RESULT: { short_name: h.short_name, task_dir: ext.task_dir, op_kind: ext.op_kind, winner_kind: cand.winner_kind, winner_backend: cand.source, target_callable: ext.target_callable || h.target_callable || '', + apply_kind: ext.apply_kind || 'python', authored_language: cand.language || '', authored_kernel_eval_dir: cand.kernel_eval_dir || '', apply_env: cand.apply_env || '', apply_flags: cand.apply_flags || '', code_patch: cand.code_patch || cand.final_patch || '', tuning_artifact: cand.tuning_artifact || '', @@ -1805,6 +1808,7 @@ while (want('kernel') && !TIME_DEADLINE_HIT && dispatched < BUDGET && (dispatche EVAL_DIR, MODEL_PATH, GPU_ID: c.gpu_id, WORKLOAD, NOISE_BAND_PCT: NOISE_BAND, E2E_REPEATS, KERNEL_RESULT: { short_name: c.short_name, task_dir: ext.task_dir, source_path_in_sglang: ext.source_path_in_sglang, target_callable: ext.target_callable, + apply_kind: ext.apply_kind || 'python', final_patch: kl.final_patch, verified_isolated_speedup: kl.final_geomean, pct_gpu_time: c.pct_gpu_time }, CURRENT_OVERLAY: curOverlay, CURRENT_FLAGS: curFlags, CURRENT_ENV: curEnv, CURRENT_THROUGHPUT: curTput, SKILL_DIR: WORKFLOW_DIR, diff --git a/e2e_workflow/knowledge/learned/INDEX.md b/e2e_workflow/knowledge/learned/INDEX.md index 196b9f2b3..15275dd3b 100644 --- a/e2e_workflow/knowledge/learned/INDEX.md +++ b/e2e_workflow/knowledge/learned/INDEX.md @@ -27,3 +27,4 @@ Confidence (a hint strength, not authority): ★ noise/unverified · ★★ sing - engagement verification: one-shot stderr banner + log grep ★★★ — (method-verify-engagement.md) - e2e A/B: pinned port, interleaved, non-overlap gate ★★★ — (method-e2e-ab-harness.md) - cuda/HIP-graph-safe integration (the #1 e2e killer) ★★★ — (method-cudagraph-safe-integration.md) +- native (.cu/.hip/.cpp/CK) apply-back: in-place incremental recompile, reversible ★★☆ — (native-apply-back.md) diff --git a/e2e_workflow/knowledge/learned/native-apply-back.md b/e2e_workflow/knowledge/learned/native-apply-back.md new file mode 100644 index 000000000..691937ff9 --- /dev/null +++ b/e2e_workflow/knowledge/learned/native-apply-back.md @@ -0,0 +1,42 @@ +--- +key: native (.cu/.hip/.cpp/CK) apply-back · any gfx · sglang/vllm +type: method +confidence: ★★☆ +effect: lets an e2e win on a COMPILED kernel actually deploy (Python overlay can't shadow a .so) +confirms: 0 +last_seen: 2026-06-29 +--- +# Deploy a COMPILED-source kernel win into the live server (in-place recompile, reversible) + +- problem: the PYTHONPATH overlay (`overlay_setup.py add-module`/`add-rebind`) can shadow a `.py` but NOT + a compiled kernel — the change only takes effect after the package's `.so`/code-object is RECOMPILED in + place. So a `.cu/.hip/.cpp/CK` winner that passes the isolated oracle is otherwise rejected + `no_rebind_seam`/`editable=false` and never reaches e2e. HL's `apply_kernel_patch` handles this but with + hardcoded framework paths + whole-package `pip install -e .`; we do it general + incremental instead. +- detection: a winner is NATIVE if `KERNEL_RESULT.apply_kind=="native"` OR the `code_patch`/`final_patch` + touches a compiled suffix `{.cu .cuh .hip .cpp .cc .cxx .c .h .hpp}`. kernel_extractor reports + `editable=true, apply_kind:"native"` for an op that resolves to a rebuildable native source (source + + a discoverable build seam shipped in the install); `editable=false` only for opaque prebuilt libs / read-only. +- apply: framework-agnostic plumbing lives in `scripts/overlay_setup.py` (`add-native` / `verify-native` / + `revert` / `gc-stale`); it NEVER invents a build command — the integrator DISCOVERS the install's own + incremental build (the same way benchmark_engineer does) and passes it via `--build-cmd`/`--invalidate-cache`. + Discover by walking UP from the source file (marker-driven, no hardcoded abs paths), first match wins: + · `config.yaml/json compile_command` / `task_runner.py compile` / `Makefile` / `build.sh` → use verbatim. + · torch `cpp_extension.load(name=…)`/`load_inline` → `--invalidate-cache $TORCH_EXTENSIONS_DIR/` + (set TORCH_EXTENSIONS_DIR per-eval + local-arch env), rebuilds on next import. + · aiter cpp_itfs `MD_NAME` driver beside source → `--invalidate-cache /_`. + · ninja `build.ninja` / CMake `build/` up-tree → build only the affected target (incremental). + · fallback: package editable rebuild — LOG it as COARSE (never silent). No seam at all → `no_build_seam`. +- verify: `verify-native` confirms the built artifact actually changed (catches a silent no-op build → + `native_build_no_op`). Only measure after FRESH_BUILD_OK. +- A/B order (native mutates the install GLOBALLY, can't toggle per-leg by env): measure REF leg FIRST on the + clean install (`OVERLAY_PYTHONPATH=$CURRENT_OVERLAY`), THEN `add-native`, THEN CAND leg (native is live in + the install; prior accepted .py overlays stay active). Mixed .py+native: `add-module` the .py into $CAND too, + CAND leg uses `OVERLAY_PYTHONPATH=$CAND`. +- reversibility (MANDATORY): wrap the CAND measurement so EVERY exit path (accept/reject/crash/timeout) runs + `overlay_setup.py revert --overlay $CAND` → install restored byte-exact; then assert `git -C status + --porcelain` empty. Runner calls `gc-stale --root $EVAL_DIR` at session start to clean a crashed run's + still-applied native overlay. Serialize: never two native A/Bs against one install concurrently (the e2e + gate leases all GPUs, so this holds by construction). +- the graph-capture-safety + memory-footprint gates apply to native kernels exactly as to authored ones. +- source: Phase 2 of the native apply-back design (DESIGN_native_apply_back.md); plumbing committed as Phase 1. diff --git a/e2e_workflow/roles/e2e_integrator.md b/e2e_workflow/roles/e2e_integrator.md index d77898ab5..5684dc5f5 100644 --- a/e2e_workflow/roles/e2e_integrator.md +++ b/e2e_workflow/roles/e2e_integrator.md @@ -59,7 +59,9 @@ Inputs: `EVAL_DIR`, `MODEL_PATH`, `BACKEND` (sglang|vllm), `GPU_ID`, `WORKLOAD`, (default 0.5), `E2E_REPEATS` (default 7; repeats per leg of the interleaved A/B), `KERNEL_RESULT` (task_dir, source_path_in_sglang, target_callable, final_patch.diff, verified_isolated_speedup, pct_gpu_time; for a HEAD-op winner also: `op_kind`, `winner_kind` -∈ {env,flag,patch}, `apply_env`, `apply_flags`, `code_patch`, `tuning_artifact`, `parity_note`), +∈ {env,flag,patch}, `apply_env`, `apply_flags`, `code_patch`, `tuning_artifact`, `parity_note`; +`apply_kind` ∈ {python (default), native} — when `native`, a `patch` is deployed via the NATIVE patch +path (in-place recompile through `overlay_setup.py add-native`) instead of a PYTHONPATH overlay), `CURRENT_OVERLAY`, `CURRENT_FLAGS`/`CURRENT_ENV`, `CURRENT_THROUGHPUT`, `SKILL_DIR`. **ACCURACY GATE (only if `ACCURACY_GATE=gsm8k` is in your inputs; else use the normal parity gate).** @@ -113,12 +115,69 @@ unchanged; you just also persist the diagnostics the deep feedback/harness-refin under `_triton_kernels/...`; shadow the file the alias actually resolves to), then `cp` that live file into `$CAND/.py`, apply the diff to the COPY, and `check` engagement via PYTHONPATH. The overlay shadows the install at import time — the live tree is never edited. - - **HARD RULE — never mutate site-packages.** The overlay is the ONLY integration mechanism; editing - `/sgl-workspace/aiter` (or any installed package) in place is forbidden (non-reversible; corrupts the - baseline leg since both legs then import the edit; contaminates other sessions on the box). Before AND - after every gate leg, assert the install is clean: `git -C /sgl-workspace/aiter status --porcelain` - (ignoring `*/flydsl_cache/`) MUST be empty. If you edited it while exploring, `git -C /sgl-workspace/aiter - checkout -- ` to restore before measuring. A win that only exists as a live-tree edit is `rejected`. + - **HARD RULE — never mutate site-packages by hand.** The Python overlay above is the ONLY integration + mechanism for `.py` changes; hand-editing `/sgl-workspace/aiter` (or any installed package) in place is + forbidden (non-reversible; corrupts the baseline leg since both legs then import the edit; contaminates + other sessions on the box). Before AND after every gate leg, assert the install is clean: + `git -C /sgl-workspace/aiter status --porcelain` (ignoring `*/flydsl_cache/`) MUST be empty. If you edited + it while exploring, `git -C /sgl-workspace/aiter checkout -- ` to restore before measuring. A win + that only exists as a hand-made live-tree edit is `rejected`. (The ONE sanctioned exception is the + **NATIVE patch path** below, which mutates the install via `overlay_setup.py add-native` — reversibly, + manifest-tracked, and always reverted in a `finally` — for compiled kernels that CANNOT be shadowed by a + PYTHONPATH overlay.) + + - **NATIVE patch (a compiled-source kernel: `.cu/.cuh/.hip/.cpp/.cc/.cxx/.c/.h/.hpp`).** A PYTHONPATH + overlay can shadow a `.py` but NOT a compiled kernel — the change only takes effect after the package's + `.so`/code-object is **recompiled** in place. Detect this case from the patch: if `code_patch`/`final_patch` + touches any compiled-suffix file, OR `KERNEL_RESULT.apply_kind=="native"`, use this path instead of + `add-module`. It is general (any framework / build system) and never invents a build command — you + DISCOVER the install's own incremental build the same way the benchmark_engineer does, then hand it to + `add-native`. Steps: + 1. **Resolve the live installed source file(s)** the diff targets (same resolution as the patch case: + `python3 -c "import ; print(.__file__)"`; honor lazy aliases). The diff applies to the LIVE + file, not a `ws` copy. + 2. **Discover the incremental build** by walking UP from the source file (marker-file driven, no + hardcoded abs paths) — first match wins: + - a `config.yaml`/`config.json` near the source declaring `compile_command`, or a `task_runner.py + compile` / `Makefile` / `build.sh` → use it verbatim (ninja/make are timestamp-incremental). + - **torch `cpp_extension.load(name=...)` / `load_inline`** in a `.py` near the source → no explicit + build cmd; pass `--invalidate-cache "$TORCH_EXTENSIONS_DIR/"` (set `TORCH_EXTENSIONS_DIR` to a + per-eval dir + local-arch env as benchmark_engineer does) so it rebuilds on next import. + - **aiter cpp_itfs**: a driver `.py` beside the source declaring `MD_NAME` → pass + `--invalidate-cache "/_"` so only that module rebuilds on next call. + - **aiter jit / setup.py+ninja / CMake `build/` / `build.ninja`** up-tree → build only the affected + target (e.g. `--build-cmd "ninja -C "` or the editable rebuild, which ninja keeps + incremental). + - **fallback**: the package's editable rebuild (e.g. `--build-cmd "python -m pip install -e . --no-build-isolation"`), + and LOG that this is a COARSE (whole-package) rebuild — never silent. If no build seam is + discoverable at all (opaque prebuilt lib, no source, read-only install) → `gate:"rejected"` reason + `no_build_seam`. + 3. **Apply reversibly + rebuild** (manifest-first, crash-safe): + ```bash + CAND="$EVAL_DIR/overlay/cand_"; cp -r "$CURRENT_OVERLAY"/. "$CAND"/ 2>/dev/null || mkdir -p "$CAND" + python3 "$SKILL_DIR/scripts/overlay_setup.py" add-native \ + --overlay "$CAND" --target "" --patch "" \ + [--artifact "" ...] [--invalidate-cache "" ...] \ + [--build-cmd ""] [--build-cwd ""] + python3 "$SKILL_DIR/scripts/overlay_setup.py" verify-native --overlay "$CAND" # FRESH_BUILD_OK or fail + ``` + If `verify-native` reports `FRESH_BUILD_FAIL` (artifact unchanged → the build was a silent no-op), + do NOT measure — fix the build command/cache target, or `gate:"rejected"` reason `native_build_no_op`. + 4. **Measure the A/B around the in-place change.** Native mutates the install GLOBALLY (it can't be + toggled per-leg by `OVERLAY_PYTHONPATH`), so order matters: measure the **REF leg FIRST** (install + still clean, `OVERLAY_PYTHONPATH="$CURRENT_OVERLAY"`), THEN `add-native`, THEN the **CAND leg** + (`OVERLAY_PYTHONPATH="$CURRENT_OVERLAY"` — the native change is already live in the install; any prior + accepted `.py` overlays stay active). For a MIXED patch (.py + native in one winner), also `add-module` + the `.py` parts into `$CAND` and run the CAND leg with `OVERLAY_PYTHONPATH="$CAND"`. + 5. **ALWAYS revert (do-no-harm + crash-safe).** Wrap the CAND measurement so that on EVERY exit path + (accept / reject / crash / timeout) you run `python3 "$SKILL_DIR/scripts/overlay_setup.py" revert + --overlay "$CAND"` to restore the install byte-exact, then assert the install is clean + (`git -C status --porcelain` empty). At session start the runner also calls + `overlay_setup.py gc-stale --root "$EVAL_DIR"` to clean any native overlay a crashed prior run left + applied. **Serialize**: never run two native A/Bs against the same install concurrently (the e2e gate + already leases all GPUs, so this holds by construction; do not parallelize native integrate legs). + The graph-capture-safety + memory-footprint gates below apply to native kernels exactly as to authored + ones. - **authored** (a from-scratch NEW implementation written by the kernel layer's author mode — there is NO installed source file to patch; instead we REBIND the op's call site to the new kernel): the authored implementation + its final patch live under diff --git a/e2e_workflow/roles/kernel_extractor.md b/e2e_workflow/roles/kernel_extractor.md index a79028b98..93423d9cd 100644 --- a/e2e_workflow/roles/kernel_extractor.md +++ b/e2e_workflow/roles/kernel_extractor.md @@ -37,6 +37,15 @@ file), `launcher_hint` (launcher seam), `bound_type`), `CURRENT_FLAGS`/`CURRENT_ `short_name` / the `module:attr` target). Confirm it's truly editable (Triton/custom/aiter) — if it resolves to a library GEMM/attention, STOP and report `editable=false` (it belongs to the Config Tuner, not here). + - **Native (compiled-source) kernels ARE editable.** If the op resolves to a COMPILED source file + (`.cu/.cuh/.hip/.cpp/.cc/.cxx/.c/.h/.hpp`) that lives in a REBUILDABLE source tree (the install ships + the source + a discoverable build seam: a `config.yaml compile_command` / `task_runner.py` / `Makefile` + / `build.sh` / `build.ninja`/CMake `build/` / a torch `cpp_extension.load(name=…)` driver / an aiter + cpp_itfs `MD_NAME` driver up-tree), report `editable=true` and `apply_kind:"native"` with + `target_callable:""` (the kernel layer authors a native diff; the e2e Integrator deploys it via the + NATIVE patch path — in-place recompile through `overlay_setup.py add-native`). Only report + `editable=false` for a truly OPAQUE prebuilt library with no shipped source (e.g. closed hipBLASLt) or a + read-only install. Default `apply_kind:"python"` for Triton/`.py` kernels (existing behavior — unchanged). 2. **Capture shapes + oracle** from a live server using `scripts/capture_shapes.py` via a temporary capture overlay, driven by the SAME workload as the profile so shapes match the regime: ```bash From 12e33ad42618a6fbd6d893d34e6363bd530592d1 Mon Sep 17 00:00:00 2001 From: zihaoan2 Date: Tue, 30 Jun 2026 05:07:49 +0000 Subject: [PATCH 3/5] feat(interface): forward optional handoff.task -> A.task steering hint General, default-off passthrough so a standalone e2e_workflow run can steer the System Architect / optimize prompts (e.g. 'only optimize compiled-source/HIP kernels; skip Triton'). Omitted => byte-identical (no hint). Co-Authored-By: Claude Opus 4.8 --- interface/run_e2e.py | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/interface/run_e2e.py b/interface/run_e2e.py index caee5b06d..81dec0375 100644 --- a/interface/run_e2e.py +++ b/interface/run_e2e.py @@ -134,6 +134,11 @@ def map_args(h: dict, timeout_s: int | None = None) -> dict: # A/B — e.g. 1 repeat per leg is enough to PROVE both legs ran). General. if h.get("e2e_repeats") is not None: ps_args["e2e_repeats"] = int(h["e2e_repeats"]) + # Optional free-text steering hint forwarded to the workflow's System Architect + # (Strategize) + optimize-phase prompts via A.task. General, default-off: e.g. + # "only optimize compiled-source/HIP kernels; skip Triton/.py". Omitted => no hint. + if h.get("task"): + ps_args["task"] = str(h["task"]) # Carried cross-phase state (the prior workflow return's `state`), so a # resume continues from where a previous phase invocation left off. if h.get("state"): From 1a3875e4e09fb1e0d7f1c00d2f5797ea9ae5bc51 Mon Sep 17 00:00:00 2001 From: zihaoan2 Date: Tue, 30 Jun 2026 08:40:36 +0000 Subject: [PATCH 4/5] Remove apply_kind axis: derive deploy mechanism from patch content The python|native exclusive axis prevented a run from composing backends across languages. Drop it everywhere (schema fields + JS threading + role docs + knowledge card). A patch's deploy mechanism is now derived per-hunk from the file suffix: compiled-suffix hunks recompile in place via overlay_setup.py add-native; .py hunks deploy as a PYTHONPATH overlay. One winner's patch may carry both -> each hunk routes by its own suffix, so a run can bank, e.g., a Triton norm and a native HIP GEMM in the same pass. Co-Authored-By: Claude Opus 4.8 --- e2e_workflow/e2e_workflow.js | 4 -- .../knowledge/learned/native-apply-back.md | 10 +-- e2e_workflow/roles/e2e_integrator.md | 65 +++++++++++++------ e2e_workflow/roles/kernel_extractor.md | 13 ++-- 4 files changed, 60 insertions(+), 32 deletions(-) diff --git a/e2e_workflow/e2e_workflow.js b/e2e_workflow/e2e_workflow.js index 05560d6f7..265fb203f 100644 --- a/e2e_workflow/e2e_workflow.js +++ b/e2e_workflow/e2e_workflow.js @@ -367,7 +367,6 @@ const EXTRACT_OP_SCHEMA = obj({ dtype: { type: 'string' }, synthesized: { type: 'boolean' }, regimes_captured: arrStr, candidate_backends: arrStr, reference_io_sha256: { type: 'string' }, target_callable: { type: 'string' }, // module:attr rebind seam for an authored kernel ('' if none) - apply_kind: { type: 'string' }, // 'python' (default) | 'native' (compiled-source kernel -> add-native in-place recompile) smoke: { type: 'string' }, notes: { type: 'string' }, }, ['op_kind', 'task_dir', 'smoke']); @@ -388,7 +387,6 @@ const EXTRACT_SCHEMA = obj({ source_path_in_sglang: { type: 'string' }, target_callable: { type: 'string' }, num_cases: { type: 'number' }, regimes_captured: arrStr, candidate_backends: arrStr, build: { type: 'boolean' }, unittest_smoke: { type: 'string' }, - apply_kind: { type: 'string' }, // 'python' (default) | 'native' reference_io_sha256: { type: 'string' }, notes: { type: 'string' }, }, ['editable', 'task_dir', 'unittest_smoke']); @@ -1603,7 +1601,6 @@ if (want('head') && headQueue.length && HEAD_BUDGET > 0) { KERNEL_RESULT: { short_name: h.short_name, task_dir: ext.task_dir, op_kind: ext.op_kind, winner_kind: cand.winner_kind, winner_backend: cand.source, target_callable: ext.target_callable || h.target_callable || '', - apply_kind: ext.apply_kind || 'python', authored_language: cand.language || '', authored_kernel_eval_dir: cand.kernel_eval_dir || '', apply_env: cand.apply_env || '', apply_flags: cand.apply_flags || '', code_patch: cand.code_patch || cand.final_patch || '', tuning_artifact: cand.tuning_artifact || '', @@ -1808,7 +1805,6 @@ while (want('kernel') && !TIME_DEADLINE_HIT && dispatched < BUDGET && (dispatche EVAL_DIR, MODEL_PATH, GPU_ID: c.gpu_id, WORKLOAD, NOISE_BAND_PCT: NOISE_BAND, E2E_REPEATS, KERNEL_RESULT: { short_name: c.short_name, task_dir: ext.task_dir, source_path_in_sglang: ext.source_path_in_sglang, target_callable: ext.target_callable, - apply_kind: ext.apply_kind || 'python', final_patch: kl.final_patch, verified_isolated_speedup: kl.final_geomean, pct_gpu_time: c.pct_gpu_time }, CURRENT_OVERLAY: curOverlay, CURRENT_FLAGS: curFlags, CURRENT_ENV: curEnv, CURRENT_THROUGHPUT: curTput, SKILL_DIR: WORKFLOW_DIR, diff --git a/e2e_workflow/knowledge/learned/native-apply-back.md b/e2e_workflow/knowledge/learned/native-apply-back.md index 691937ff9..c3fabb56d 100644 --- a/e2e_workflow/knowledge/learned/native-apply-back.md +++ b/e2e_workflow/knowledge/learned/native-apply-back.md @@ -13,10 +13,12 @@ last_seen: 2026-06-29 place. So a `.cu/.hip/.cpp/CK` winner that passes the isolated oracle is otherwise rejected `no_rebind_seam`/`editable=false` and never reaches e2e. HL's `apply_kernel_patch` handles this but with hardcoded framework paths + whole-package `pip install -e .`; we do it general + incremental instead. -- detection: a winner is NATIVE if `KERNEL_RESULT.apply_kind=="native"` OR the `code_patch`/`final_patch` - touches a compiled suffix `{.cu .cuh .hip .cpp .cc .cxx .c .h .hpp}`. kernel_extractor reports - `editable=true, apply_kind:"native"` for an op that resolves to a rebuildable native source (source + - a discoverable build seam shipped in the install); `editable=false` only for opaque prebuilt libs / read-only. +- detection: there is NO `apply_kind` axis — the deploy mechanism is derived per-hunk from the patch + content. A hunk is NATIVE if it touches a compiled suffix `{.cu .cuh .hip .cpp .cc .cxx .c .h .hpp}`; + a `.py` hunk deploys as a PYTHONPATH overlay. One winner's patch may carry both → each hunk routes by + its own suffix; the modes compose, they are not exclusive. kernel_extractor reports `editable=true` + (with `target_language` hip/ck) for an op that resolves to a rebuildable native source (source + a + discoverable build seam shipped in the install); `editable=false` only for opaque prebuilt libs / read-only. - apply: framework-agnostic plumbing lives in `scripts/overlay_setup.py` (`add-native` / `verify-native` / `revert` / `gc-stale`); it NEVER invents a build command — the integrator DISCOVERS the install's own incremental build (the same way benchmark_engineer does) and passes it via `--build-cmd`/`--invalidate-cache`. diff --git a/e2e_workflow/roles/e2e_integrator.md b/e2e_workflow/roles/e2e_integrator.md index 5684dc5f5..1880216e6 100644 --- a/e2e_workflow/roles/e2e_integrator.md +++ b/e2e_workflow/roles/e2e_integrator.md @@ -59,11 +59,15 @@ Inputs: `EVAL_DIR`, `MODEL_PATH`, `BACKEND` (sglang|vllm), `GPU_ID`, `WORKLOAD`, (default 0.5), `E2E_REPEATS` (default 7; repeats per leg of the interleaved A/B), `KERNEL_RESULT` (task_dir, source_path_in_sglang, target_callable, final_patch.diff, verified_isolated_speedup, pct_gpu_time; for a HEAD-op winner also: `op_kind`, `winner_kind` -∈ {env,flag,patch}, `apply_env`, `apply_flags`, `code_patch`, `tuning_artifact`, `parity_note`; -`apply_kind` ∈ {python (default), native} — when `native`, a `patch` is deployed via the NATIVE patch -path (in-place recompile through `overlay_setup.py add-native`) instead of a PYTHONPATH overlay), +∈ {env,flag,patch}, `apply_env`, `apply_flags`, `code_patch`, `tuning_artifact`, `parity_note`), `CURRENT_OVERLAY`, `CURRENT_FLAGS`/`CURRENT_ENV`, `CURRENT_THROUGHPUT`, `SKILL_DIR`. +There is no `apply_kind` axis. The deploy mechanism for a `patch` is derived from the patch +content itself: a hunk touching a `.py` file deploys as a PYTHONPATH overlay (reversible, add-module); +a hunk touching a compiled-source file (`.cu/.cuh/.hip/.cpp/.cc/.cxx/.c/.cl`) deploys via the NATIVE +path (in-place recompile through `overlay_setup.py add-native`). One winner's patch may contain both, +in which case each hunk routes by its own suffix — the modes compose, they are not exclusive. + **ACCURACY GATE (only if `ACCURACY_GATE=gsm8k` is in your inputs; else use the normal parity gate).** For a QUANTIZED kernel, byte-exact greedy parity is the WRONG bar (a within-tolerance kernel rounds differently → flips borderline argmaxes → over-rejects valid kernels). Instead, score TASK ACCURACY: @@ -128,8 +132,8 @@ unchanged; you just also persist the diagnostics the deep feedback/harness-refin - **NATIVE patch (a compiled-source kernel: `.cu/.cuh/.hip/.cpp/.cc/.cxx/.c/.h/.hpp`).** A PYTHONPATH overlay can shadow a `.py` but NOT a compiled kernel — the change only takes effect after the package's - `.so`/code-object is **recompiled** in place. Detect this case from the patch: if `code_patch`/`final_patch` - touches any compiled-suffix file, OR `KERNEL_RESULT.apply_kind=="native"`, use this path instead of + `.so`/code-object is **recompiled** in place. Detect this case from the patch alone: if a + `code_patch`/`final_patch` hunk touches any compiled-suffix file, that hunk uses this path instead of `add-module`. It is general (any framework / build system) and never invents a build command — you DISCOVER the install's own incremental build the same way the benchmark_engineer does, then hand it to `add-native`. Steps: @@ -169,15 +173,29 @@ unchanged; you just also persist the diagnostics the deep feedback/harness-refin (`OVERLAY_PYTHONPATH="$CURRENT_OVERLAY"` — the native change is already live in the install; any prior accepted `.py` overlays stay active). For a MIXED patch (.py + native in one winner), also `add-module` the `.py` parts into `$CAND` and run the CAND leg with `OVERLAY_PYTHONPATH="$CAND"`. - 5. **ALWAYS revert (do-no-harm + crash-safe).** Wrap the CAND measurement so that on EVERY exit path - (accept / reject / crash / timeout) you run `python3 "$SKILL_DIR/scripts/overlay_setup.py" revert - --overlay "$CAND"` to restore the install byte-exact, then assert the install is clean - (`git -C status --porcelain` empty). At session start the runner also calls - `overlay_setup.py gc-stale --root "$EVAL_DIR"` to clean any native overlay a crashed prior run left - applied. **Serialize**: never run two native A/Bs against the same install concurrently (the e2e gate - already leases all GPUs, so this holds by construction; do not parallelize native integrate legs). - The graph-capture-safety + memory-footprint gates below apply to native kernels exactly as to authored - ones. + 5. **Native wins COMPOUND — keep on accept, revert only on reject/crash.** Native is just another + accepted-overlay layer; it must stack with prior wins (Triton/`.py` overlays AND earlier native + wins) exactly like the PYTHONPATH overlay does — do NOT revert an accepted native change between + kernels, or the run can only ever bank one native kernel. Concretely, keep TWO native overlays: + - a **cumulative** `$EVAL_DIR/overlay/native_accepted` (all ACCEPTED native changes; its manifest + accumulates) — this stays APPLIED in the install across the whole run, so every later REF leg + already includes it and later CAND legs stack on top. + - a **candidate** `$CAND_NATIVE` for the change under test (apply with `add-native --overlay + "$CAND_NATIVE"`). + On the gate outcome: + - **reject / crash / timeout** → `revert --overlay "$CAND_NATIVE"` (restore byte-exact), assert the + install matches the cumulative state (only `native_accepted` changes remain). + - **accept** → DO NOT revert; fold the candidate into `native_accepted` (re-apply it against that + overlay so its manifest records the source+artifact backups), leaving it live for subsequent + kernels. Carry `native_accepted` forward alongside `CURRENT_OVERLAY` (the `.py` overlay), just like + accepted flags/env. + Crash-safety: every applied change is in a manifest before it mutates the install, so + `overlay_setup.py gc-stale --root "$EVAL_DIR"` (run at session start) restores a crashed run's + install. The install is only fully reverted to pristine at **Finalize**, AFTER the bundle has + captured the rebuilt artifacts + a re-apply recipe (do-no-harm to the box). **Serialize**: never run + two native A/Bs against the same install concurrently (the e2e gate already leases all GPUs, so this + holds; do not parallelize native integrate legs). The graph-capture-safety + memory-footprint gates + below apply to native kernels exactly as to authored ones. - **authored** (a from-scratch NEW implementation written by the kernel layer's author mode — there is NO installed source file to patch; instead we REBIND the op's call site to the new kernel): the authored implementation + its final patch live under @@ -338,12 +356,21 @@ Return JSON: Inputs: `EVAL_DIR`, the final accepted overlay, accepted config (flags/env), all accepted kernel patches, `BASELINE_THROUGHPUT`, `SKILL_DIR`. -1. Assemble the deliverable bundle in `EVAL_DIR/final/`: the accepted overlay dir, a concatenated - `final_patch.diff` (all accepted kernel patches), and a `final_launch.sh` that reproduces the - optimized server (sets `BACKEND=`, `PYTHONPATH=`, the accepted flags/env, and runs - the bench via bench_e2e.sh + its adapter). This is the spec deliverable: "complete patch + launch/benchmark script". +1. Assemble the deliverable bundle in `EVAL_DIR/final/`: the accepted Python overlay dir, a concatenated + `final_patch.diff` (all accepted kernel patches — `.py` AND native), and a `final_launch.sh` that + reproduces the optimized server. The bundle must carry BOTH modes the run banked: + - **Python/Triton wins** → `PYTHONPATH=` in `final_launch.sh` (as before). + - **Native (.cu/.hip/CK) wins** → copy the cumulative `native_accepted` overlay (its manifest + source + backups + rebuilt artifacts) into `EVAL_DIR/final/native/`, and have `final_launch.sh` **re-apply + native before serving** (`overlay_setup.py add-native ... --build-cmd ` per recorded entry, + OR drop the captured rebuilt artifacts into the install) so a fresh checkout reproduces the compiled + change. Both modes compose — a run that banked, say, a Triton norm + a native GEMM ships both. 2. Do a final warm-server bench of the assembled bundle to confirm the combined result matches the - sum of accepted milestones (combined effects can interact). Record it. + sum of accepted milestones (Python + native effects can interact). Record it. +3. **Restore the box** (do-no-harm): after the bundle has captured the native artifacts/recipe, revert the + install to pristine — `overlay_setup.py revert --overlay "$EVAL_DIR/overlay/native_accepted"` — and + assert `git -C status --porcelain` is empty. The deliverable bundle (not a mutated site-packages) + is the artifact; the user re-applies via `final_launch.sh`. Return JSON: ```json diff --git a/e2e_workflow/roles/kernel_extractor.md b/e2e_workflow/roles/kernel_extractor.md index 93423d9cd..13467f2b9 100644 --- a/e2e_workflow/roles/kernel_extractor.md +++ b/e2e_workflow/roles/kernel_extractor.md @@ -41,11 +41,14 @@ file), `launcher_hint` (launcher seam), `bound_type`), `CURRENT_FLAGS`/`CURRENT_ (`.cu/.cuh/.hip/.cpp/.cc/.cxx/.c/.h/.hpp`) that lives in a REBUILDABLE source tree (the install ships the source + a discoverable build seam: a `config.yaml compile_command` / `task_runner.py` / `Makefile` / `build.sh` / `build.ninja`/CMake `build/` / a torch `cpp_extension.load(name=…)` driver / an aiter - cpp_itfs `MD_NAME` driver up-tree), report `editable=true` and `apply_kind:"native"` with - `target_callable:""` (the kernel layer authors a native diff; the e2e Integrator deploys it via the - NATIVE patch path — in-place recompile through `overlay_setup.py add-native`). Only report - `editable=false` for a truly OPAQUE prebuilt library with no shipped source (e.g. closed hipBLASLt) or a - read-only install. Default `apply_kind:"python"` for Triton/`.py` kernels (existing behavior — unchanged). + cpp_itfs `MD_NAME` driver up-tree), report `editable=true` with `target_callable:""` and route to the + kernel layer with `target_language` ∈ {hip, ck} (the kernel layer authors a native diff against the + compiled source). The deploy mechanism is NOT decided here — the e2e Integrator derives it from the + resulting patch: any hunk touching a compiled-suffix file is deployed via the NATIVE patch path + (in-place recompile through `overlay_setup.py add-native`); `.py` hunks deploy as a PYTHONPATH overlay. + Only report `editable=false` for a truly OPAQUE prebuilt library with no shipped source (e.g. closed + hipBLASLt) or a read-only install. Triton/`.py` kernels remain `editable=true` with their usual + `target_language` (existing behavior — unchanged). 2. **Capture shapes + oracle** from a live server using `scripts/capture_shapes.py` via a temporary capture overlay, driven by the SAME workload as the profile so shapes match the regime: ```bash From b6bb644de47f920c9c90650eba0355993e9179ca Mon Sep 17 00:00:00 2001 From: zihaoan2 Date: Tue, 30 Jun 2026 13:56:04 +0000 Subject: [PATCH 5/5] extractor: resolve the REAL live kernel + dispatch seam (no guessing, no hardcode) A head GEMM authored a real 1.92x isolated kernel on Qwen3-4B/vLLM that engaged 0 times e2e and was gate-rejected. Root cause was 100% seam resolution: the extractor returned a guessed launcher (aiter.tuned_gemm:gemm_a16w16) inferred from the trace kernel name, but the live dense GEMM dispatched to closed hipBLASLt (the Cijk_* Tensile kernels) -- a different module, a shape-gated branch the model never hit, a closed leaf. The authored kernel bound a symbol nothing calls. Make seam resolution evidence-driven and framework-agnostic: - New method card knowledge/learned/method-resolve-live-seam.md: classify the profiled GPU kernel's backend family from the trace + owning .so; walk the LIVE installed dispatch source under the server's env + captured shapes; CONFIRM with a temporary throwaway [seam-probe] passthrough before authoring; closed-library leaf (Tensile/hipBLASLt) -> editable=false -> Config Tuner, not author; bind the nearest common chokepoint when decode/prefill branches diverge; prove engagement by re-observation. - kernel_extractor.md: replace the hardcoded "dense GEMM seam is aiter.tuned_gemm:gemm_a16w16" guidance (both PHASE=extract and PHASE=extract_op) with the resolve-then-confirm procedure; a Cijk_*/ Tensile/hipBLASLt leaf is closed -> editable=false. - gemm_attention_backends.md: scope the aiter seam to an EXAMPLE of one resolved path (sglang/gfx942), not the default; require resolution first. - INDEX.md: register the new method card. No backend is privileged; no module name is hardcoded; the seam is whatever the live dispatch actually calls for the captured shapes under the live env. Co-Authored-By: Claude Opus 4.8 --- .../knowledge/gemm_attention_backends.md | 24 +++-- e2e_workflow/knowledge/learned/INDEX.md | 1 + .../learned/method-resolve-live-seam.md | 101 ++++++++++++++++++ e2e_workflow/roles/kernel_extractor.md | 39 +++++-- 4 files changed, 150 insertions(+), 15 deletions(-) create mode 100644 e2e_workflow/knowledge/learned/method-resolve-live-seam.md diff --git a/e2e_workflow/knowledge/gemm_attention_backends.md b/e2e_workflow/knowledge/gemm_attention_backends.md index 966c31fe9..5ae6430ad 100644 --- a/e2e_workflow/knowledge/gemm_attention_backends.md +++ b/e2e_workflow/knowledge/gemm_attention_backends.md @@ -10,12 +10,24 @@ backends) **the kernel code itself** — cheapest first. > Why this exists: e2e is Amdahl-dominated. A GEMM at ~78% of GPU time only needs **1.15x** to give > ~+10% e2e. That is a far better bet than a 1.3x on a 2% kernel. Spend budget on the head first. -> **On sglang/gfx942, dense-GEMM Tier-B = aiter's per-shape DB; Tier-C = an authored Triton GEMM.** -> The live dense-GEMM path is aiter `tuned_gemm.py` (seam `aiter.tuned_gemm:gemm_a16w16`). Tune it via -> `AITER_TUNE_GEMM=1` capture → `gradlib/gemm_tuner.py` → deploy `AITER_CONFIG_GEMM_BF16`, and verify -> engagement with `AITER_LOG_TUNED_CONFIG=1`. Full recipe: **`gemm_tuning/aiter_gemm_tuning.md`**. (TunableOp / -> `HIPBLASLT_TUNING_FILE` hook the PyTorch dispatch, which this live path does not use — so they are not -> the GEMM lever here; tune aiter / author Triton instead.) +> **RESOLVE the live dense-GEMM path before picking a lever — it differs by framework/gfx/model and +> must NOT be assumed.** Follow `learned/method-resolve-live-seam.md`: classify the profiled GEMM +> kernel's backend family from the trace + owning `.so`, then walk the live dispatch source under the +> server's env + captured shapes, and confirm with a throwaway `[seam-probe]`. Only then choose the lever. +> - **Closed library leaf** (`Cijk_*`/Tensile → hipBLASLt/rocBLAS; the common large-M dense case) → +> Tier-B tuning ONLY; there is no authorable source seam. Tune the actual library: hipBLASLt solution +> sweep / `HIPBLASLT_TUNING_FILE` / TunableOp **if** the live path goes through the PyTorch dispatch, +> else the backend's own tuner. +> - **aiter `tuned_gemm.py` leaf** (observed on sglang/gfx942; seam `aiter.tuned_gemm:gemm_a16w16`) → +> Tier-B = `AITER_TUNE_GEMM=1` capture → `gradlib/gemm_tuner.py` → deploy `AITER_CONFIG_GEMM_BF16`, +> verify with `AITER_LOG_TUNED_CONFIG=1` (recipe: `gemm_tuning/aiter_gemm_tuning.md`); Tier-C = author +> a Triton GEMM bound to that confirmed seam. **This is an EXAMPLE of one resolved path, not the +> default** — on vLLM the dense GEMM dispatches through `torch.ops.vllm.rocm_unquantized_gemm` whose +> branches may be hipBLASLt (closed), `aiter.ops.triton.gemm_a16w16` (a DIFFERENT module, shape-gated), +> or skinny C++ ops; resolve which one fires for YOUR shapes. +> - **Editable Triton/CK leaf** → Tier-C authored rewrite bound to the confirmed seam (or Tier-B +> autotune within it). A flag that re-routes a closed op to an editable backend is a Tier-A/CONFIG +> lever — but only if its predicate accepts this model's shapes (verify, don't assume). ## The ladder (cheapest-first; each rung gated by the immutable oracle + e2e Amdahl + parity) diff --git a/e2e_workflow/knowledge/learned/INDEX.md b/e2e_workflow/knowledge/learned/INDEX.md index 15275dd3b..4abc5dab2 100644 --- a/e2e_workflow/knowledge/learned/INDEX.md +++ b/e2e_workflow/knowledge/learned/INDEX.md @@ -24,6 +24,7 @@ Confidence (a hint strength, not authority): ★ noise/unverified · ★★ sing - [gfx942 · prefill-dominated hybrid] stack-and-compound cluster; Amdahl pre-dispatch screen ★★★ — (editable-triton-cluster-amdahl.md) ## method (cross-model, applies to any run) +- resolve live kernel + dispatch seam (no guessing/hardcode; classify trace→walk live source→throwaway probe; closed-lib→tune not author) ★★★ — (method-resolve-live-seam.md) - engagement verification: one-shot stderr banner + log grep ★★★ — (method-verify-engagement.md) - e2e A/B: pinned port, interleaved, non-overlap gate ★★★ — (method-e2e-ab-harness.md) - cuda/HIP-graph-safe integration (the #1 e2e killer) ★★★ — (method-cudagraph-safe-integration.md) diff --git a/e2e_workflow/knowledge/learned/method-resolve-live-seam.md b/e2e_workflow/knowledge/learned/method-resolve-live-seam.md new file mode 100644 index 000000000..68fc485dc --- /dev/null +++ b/e2e_workflow/knowledge/learned/method-resolve-live-seam.md @@ -0,0 +1,101 @@ +--- +key: live-kernel + dispatch-seam resolution · any gfx · any framework · any backend +type: method +confidence: ★★★ +effect: the kernel you optimize is the kernel the GPU actually ran, bound to the seam the live path actually calls — instead of a plausible-looking guess that engages 0 times +confirms: 1 +last_seen: 2026-06-29 +--- +# Resolve the REAL live kernel + its dispatch seam (never guess the launcher from a kernel name) + +**Why this exists.** A head GEMM authored a real isolated 1.92× kernel on Qwen3-4B/vLLM, then engaged +**0 times** e2e and was gate-rejected. Root cause was 100% seam-resolution, not the apply mechanism: +the extractor returned a *guessed* launcher (`aiter.tuned_gemm:gemm_a16w16`) from the kernel's trace +name. But on that box the live dense GEMM (a) imports a **different module** if it ever went Triton +(`aiter.ops.triton.gemm_a16w16`), (b) that branch is gated by a **hardcoded shape whitelist that the +model didn't match**, and (c) the shapes actually dispatched to **closed hipBLASLt** (the `Cijk_*` +Tensile kernels in the trace). The authored kernel bound a symbol nothing calls → 0 engagement. + +**The rule: the seam is whatever the LIVE dispatch actually calls for the CAPTURED shapes under the +LIVE env — discovered by observation + reading the installed source, never assumed from the op type, +the model, the gfx, or a `launcher_hint`.** No backend (aiter/triton/ck/…) is privileged. No module +name is hardcoded. Do this BEFORE authoring, so a closed-library op is routed to tuning instead of +wasting an author cycle. + +## The two ends you must connect +- **GPU end (ground truth)**: the profiler already recorded the exact kernel that ran. That is + authoritative — the kernel that ran is whatever the trace says, not what we expect. +- **Host/Python end (the seam)**: the callable whose body decided to launch that GPU kernel. This is + what the Integrator rebinds (`.py` overlay) or patches (native recompile). Resolve THIS, don't guess it. + +## Step 1 — Classify the GPU kernel's backend FAMILY from the trace (pattern, not a fixed list) +Take the kernel name from `profile_topN.json`; if it's a mangled C++ symbol, `c++filt` it; note which +`.so`/module rocprof attributes it to (the profiler gives the owning module path per kernel — use it). +Map signature → owning backend by what the name/.so reveals, e.g.: +- `Cijk_*`, `Cgemm_*`, Tensile-mangled, or attributed to `librocblas`/`libhipblaslt` → **rocBLAS/hipBLASLt + (Tensile) = CLOSED**. No shipped source, no python seam → not authorable. +- `*wvSplitK*`, `LLMM*`, other names in an `aiter`/`vllm` `.so` → compiled C++ op (editable only if the + source ships in a rebuildable tree — see [[native-apply-back]]). +- `triton_*`, `*_fwd_kernel*`, inductor `triton_poi_/red_fused_*` → **Triton (editable `.py`)**. +- `ck_*`, composable_kernel device names → **CK** (editable C++ if source shipped). +This is reasoning from the symbol, not memorizing — any new name: demangle it, find its `.so`, decide +closed-lib vs editable-source by whether source ships up-tree. + +## Step 2 — Map that GPU kernel back to its launching Python frame (the seam), by evidence +Prefer (a); fall back to (b). Plug in REAL values — never evaluate a branch in your head from defaults. +- **(a) Profiler call-stack / API correlation.** If the trace carries host stacks (torch profiler + `with_stack=True`, or rocprofv3 HIP-API + correlation IDs), read the launching frame directly — that + function IS the seam. Cheapest, most reliable when available. +- **(b) Source walk from the op entry.** Open the **live installed** source at the op family's entry + (the registered custom op / `Linear.forward` / the attention-backend forward — found by `import pkg; + os.path.dirname` then grep, NOT from memory) and follow the ACTUAL branch tree, evaluating each + predicate against the **server's live env vars/flags** and the **captured shapes**. The branch that + fires for those shapes terminates at the real leaf callable = the seam. Read the code; substitute the + live values; do not assume which branch wins. + +## Step 3 — Confirm the seam with a TEMPORARY throwaway probe (allowed, encouraged) +A disposable one-shot python probe is the cheapest proof — it is a verification tool, NOT part of the +integration. Put it in a fresh `_seam_probe/` overlay, delete after: +- `setattr` a passthrough wrapper on the resolved `module:attr` that prints a one-time + `[seam-probe] HIT shape=<...>` to stderr, then calls the original. (For a leaf imported + *inside* a function via `from M import f` each call, patch `M.f` — the per-call import re-fetches it.) +- Run the SAME workload for a few seconds (REPEATS=0 warmup window) and grep the server log: + - **≥1 HIT per worker** → seam confirmed; record the exact `module:attr`. Proceed to author. + - **0 HITs but a concurrent re-profile still shows the target GPU kernel** → wrong seam. Either you + patched the wrong module (Step 1's family was off) or the live path bypasses python entirely + (closed lib) → go to Step 4 "closed". Re-walk Step 2; do NOT author against an unconfirmed seam. +The probe is generic: it does not name any backend, only the `module:attr` you resolved this run. + +## Step 4 — Editability is decided by the RESOLVED LEAF, not the op kind +- Leaf = **closed library** (Tensile/hipBLASLt/rocBLAS, or an opaque `.so` with no shipped source) → + `editable=false` → route to **Config Tuner** (per-shape tuning DB, or a flag/backend swap), do NOT + author. This is the case that wasted the 1.92× cycle — catch it here, pre-author. +- Leaf = **editable Triton `.py`** or **compiled source in a rebuildable tree** → `editable=true`; + record the EXACT confirmed `module:attr` (+ source file for native). Author. +- Leaf is closed NOW but an **editable alternative exists behind a flag** (an env that re-routes the + SAME op to a Triton/CK path) → this is a **CONFIG lever first**: record the flag to flip AND the seam + it then exposes. Flip → re-confirm via Step 3 against the new live path → then author. (Flipping a + flag that re-routes to editable code is the general way to "reach" an otherwise-closed op — but it + only counts if the flag's predicate actually accepts the model's shapes; verify, don't assume.) + +## Step 5 — When per-shape branches diverge, bind the nearest common chokepoint +Decode (skinny-M) and prefill (large-M) can take DIFFERENT branches → DIFFERENT leaves → maybe one +editable and one closed. If you bind a single leaf you cover only part of the shape space (and may +regress the other). Prefer the **nearest common chokepoint every shape passes through** (the op entry / +registered custom op / `Linear.forward`), patched to dispatch to the authored kernel for the target +shapes and **fall through to the original for the rest**. Confirm with Step 3 that the chokepoint sees +both regimes' shapes. + +## Step 6 — Engagement is proven by re-OBSERVATION, not an e2e wiggle +This is the post-author gate (see [[method-verify-engagement]]). After integrating, prove the OLD GPU +kernel disappeared and the NEW one appears in a re-profile (or the candidate's own one-shot banner fires +INSIDE the cudagraph-captured region — see [[method-cudagraph-safe-integration]]). A throughput delta +with no kernel-swap evidence is noise — reject it (that is exactly what caught the Qwen3-4B false win). + +## Anti-hardcode checklist (the failure mode this card prevents) +- [ ] Did NOT assume a backend from the model/gfx — classified from THIS trace's kernel names + `.so`. +- [ ] Did NOT trust a `launcher_hint`/learned-card seam blindly — confirmed with a live probe or re-profile. +- [ ] Resolved the seam by reading the **installed** source + live env + captured shapes, not from memory. +- [ ] If the leaf is a closed lib → routed to tuning, did NOT author. +- [ ] Confirmed the seam sees the shapes BEFORE authoring; confirmed kernel-swap AFTER integrating. +- source: exp/.../e2e_Qwen3-4B_v1 dense-GEMM 0-engagement reject (2026-06-29); vllm rocm_unquantized_gemm dispatch tree diff --git a/e2e_workflow/roles/kernel_extractor.md b/e2e_workflow/roles/kernel_extractor.md index 13467f2b9..82379d262 100644 --- a/e2e_workflow/roles/kernel_extractor.md +++ b/e2e_workflow/roles/kernel_extractor.md @@ -36,7 +36,11 @@ file), `launcher_hint` (launcher seam), `bound_type`), `CURRENT_FLAGS`/`CURRENT_ (`python3 -c "import sglang,os;print(os.path.dirname(sglang.__file__))"`, then grep the `short_name` / the `module:attr` target). Confirm it's truly editable (Triton/custom/aiter) — if it resolves to a library GEMM/attention, STOP and report `editable=false` (it belongs to the - Config Tuner, not here). + Config Tuner, not here). **Editability is decided by the kernel the profile shows ACTUALLY ran on the + live path — not by the op type or a hint** — resolve + confirm the live seam per + `SKILL_DIR/knowledge/learned/method-resolve-live-seam.md` (classify the GPU kernel's backend family + from the trace, walk the live dispatch source under the server's env/shapes, confirm with a temporary + throwaway `[seam-probe]` passthrough). A `Cijk_*`/Tensile/hipBLASLt leaf is CLOSED → `editable=false`. - **Native (compiled-source) kernels ARE editable.** If the op resolves to a COMPILED source file (`.cu/.cuh/.hip/.cpp/.cc/.cxx/.c/.h/.hpp`) that lives in a REBUILDABLE source tree (the install ships the source + a discoverable build seam: a `config.yaml compile_command` / `task_runner.py` / `Makefile` @@ -193,14 +197,31 @@ force real compact-operand compute: 5. Finalize `meta.json` with the `reference_io_sha256` (when an oracle file exists) and smoke-test `op_bench.py --task --backends hipblaslt --repeats 5` (gemm) so the harness is proven before the bake-off. -6. **Report a `target_callable` rebind seam** (`module:attr`) — this is where the e2e Integrator rebinds - the op's call site to an AUTHORED kernel. **For dense GEMM on sglang/gfx942 there IS a clean seam: - the live path goes through aiter's `aiter.tuned_gemm:gemm_a16w16` (and `aiter.tuned_gemm.tgemm.mm`), - not raw `F.linear`** — so return `target_callable="aiter.tuned_gemm:gemm_a16w16"` (or the specific - sglang Linear method that calls it, whichever the Integrator can monkeypatch cleanly). Confirm by - grepping the server for `tuned_gemm`/`gemm_a16w16` on the live path. For attention, the seam is the - backend forward you captured. Only return `target_callable=""` if no Python seam genuinely exists - (then an authored kernel can't be wired and a direct_light env winner still applies). +6. **Report a `target_callable` rebind seam** (`module:attr`) — where the e2e Integrator rebinds the + op's call site to an AUTHORED kernel. **Resolve this seam, never guess it** — follow + `SKILL_DIR/knowledge/learned/method-resolve-live-seam.md` (the full framework-agnostic procedure). + In short: (1) classify the profiled GPU kernel's backend FAMILY from its trace name + owning `.so` + (Tensile/`Cijk_*`/`librocblas`/`libhipblaslt` = CLOSED; `triton_*`/`*_fwd_kernel` = editable `.py`; + `ck_*` = CK; compiled `.so` op = native-if-source-ships); (2) map it back to the launching Python + frame by reading the **live installed** dispatch source (the registered op / Linear.forward / the + attention forward — resolved by `import pkg; os.path.dirname` + grep, NOT from memory), evaluating + each branch against the **server's live env/flags** and the **captured shapes**; (3) CONFIRM with a + **temporary throwaway probe** — `setattr` a one-shot `[seam-probe] HIT` passthrough on + the resolved callable, run the same workload a few seconds, grep the log: ≥1 hit/worker → confirmed; + 0 hits but the GPU kernel still appears → wrong seam, re-resolve (do NOT report it). The probe names + no backend and is deleted after. + - **If the resolved leaf is a CLOSED library** (Tensile/hipBLASLt/rocBLAS — the common dense-GEMM + case; the `Cijk_*` kernels are hipBLASLt) → there is NO authorable python/source seam. Report + `editable=false`, `target_callable=""` → it belongs to the **Config Tuner** (per-shape tuning DB, + or a flag/backend swap), not the kernel author. Do NOT return a plausible aiter/triton symbol you + did not confirm fires — that is exactly the 0-engagement reject this guards against. + - **If a flag re-routes the same op to an editable backend**, report it as a CONFIG lever (the flag + to flip) plus the seam it then exposes — but only after confirming (step 3) the flag's predicate + actually accepts THIS model's shapes; many such flags are shape-whitelisted. + - When decode and prefill take different branches/leaves, prefer the **nearest common chokepoint** + (op entry) the Integrator can patch with shape-gated fall-through, not a single leaf. For attention, + the seam is the backend forward you captured (confirm with the same probe). Only return + `target_callable=""` when no python seam genuinely exists. > **Shapes must be the REAL ones the server issues — and they MUST span BOTH regimes.** A head GEMM > serves many M buckets: the **decode** regime at small M = the steady-state running batch (M ≈ `WORKLOAD.conc`,