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 196b9f2b3..4abc5dab2 100644 --- a/e2e_workflow/knowledge/learned/INDEX.md +++ b/e2e_workflow/knowledge/learned/INDEX.md @@ -24,6 +24,8 @@ 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) +- native (.cu/.hip/.cpp/CK) apply-back: in-place incremental recompile, reversible ★★☆ — (native-apply-back.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/knowledge/learned/native-apply-back.md b/e2e_workflow/knowledge/learned/native-apply-back.md new file mode 100644 index 000000000..c3fabb56d --- /dev/null +++ b/e2e_workflow/knowledge/learned/native-apply-back.md @@ -0,0 +1,44 @@ +--- +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: 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`. + 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..1880216e6 100644 --- a/e2e_workflow/roles/e2e_integrator.md +++ b/e2e_workflow/roles/e2e_integrator.md @@ -62,6 +62,12 @@ verified_isolated_speedup, pct_gpu_time; for a HEAD-op winner also: `op_kind`, ` ∈ {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: @@ -113,12 +119,83 @@ 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 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: + 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. **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 @@ -279,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 a79028b98..82379d262 100644 --- a/e2e_workflow/roles/kernel_extractor.md +++ b/e2e_workflow/roles/kernel_extractor.md @@ -36,7 +36,23 @@ 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` + / `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` 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 @@ -181,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`, 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) 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"):