feat(e2e): native (.cu/.hip/.cpp/CK) apply-back — in-place incremental recompile, reversible

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)
 

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)

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] <module:attr> 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

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/<name>`
+    (set TORCH_EXTENSIONS_DIR per-eval + local-arch env), rebuilds on next import.
+  · aiter cpp_itfs `MD_NAME` driver beside source → `--invalidate-cache <cpp_itfs_root>/<md_name>_<hash>`.
+  · 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 <pkg> 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.