mpk-faithful-gate

# MPK Faithful Per-Task Gate ## Why this exists (read first — it's the whole point) Two tempting measures LIE for the DSv3 decode campaign: - A **standalone green-ctx kernel bench** (ferret/cpp_examples) runs the kernel alone on a fresh CUDA context. It MIS-RANKS — a kernel that wins standalone can be NULL or a REGRESS once it's compiled into the shared megakernel at the production grid with co-resident tasks (3 confirmed "standalone-doesn't-transfer" cases: kv-up CUDA-core, o_proj split-K, dense fine-N TP8). - The **whole-megakernel e2e** per-token latency hides which task moved. The trusted measure is the **FAITHFUL in-MPK per-task** number: compile the kernel INTO the real megakernel, launch the PRODUCTION persistent grid (`grid.x = num_workers`, 136 on a B200), turn the persistent profiler ON, and extract THIS task's per-instance span: - **slowCTA** = max-over-CTAs of (end − begin) = the slowest single-CTA BODY (per-instance compute). THIS is the verdict metric. Not P50, not per-kernel-aggregate, not wall. - **wall** = max(end) − min(begin) = the isolated makespan (includes dispatch stagger). - **sumCTA** = total busy CTA-time = the ANTI-SLOUGHING guard. A candidate that lowers slowCTA by spreading the SAME work onto more CTAs (no real speedup) shows a RISING sumCTA — reject it. Judge on slowCTA AND wall AND sumCTA together. - **cos** ≥ 0.99 vs a PyTorch reference (correctness; never trade it for speed). ## The existing gates — your templates (copy structure, don't reinvent) | Family | Gate file | Notes | |---|---|---| | Dense FP8 GEMM (qkv_a/q_b/q_b_pe/kv_b/o_proj/shared_*/router) | `tests/runtime_python/blackwell/sm100_fp8_gemm_dense/test_fp8_gemm_dense_*_pk_testmode.py` + `_build_helper.py` (`SHAPE_REGISTRY`) + `_faithful_helper.py` | M=1 decode; shapes are named in `SHAPE_REGISTRY` (add a row for a new dense shape) | | Routed group-GEMM W13/W2 | `tests/runtime_python/blackwell/sm100_fp8_group_gemm_decode/test_fp8_group_gemm_faithful_pertask.py` | decode geometry via `MPK_GG_DECODE_M1`; active experts via `meta` row-1 mask | | (template core, kernel-agnostic) | `sm100_fp8_gemm_dense/_faithful_helper.py` → `profiled_per_task_latency_by_id`, `resolve_num_workers`, `make_profiler_tensor` | import these; do not rewrite | | Candidate-overlay bridge | `~/ferret/scripts/faithful_eval.py` (KIND map + `_run_test` + `_run_group_gemm`) | drives baseline vs shadow-overlay candidate in two subprocesses | | GPU broker | `~/gpu_broker/` (`gpu_gate.sh`, `gpu_pool.conf`, `acquire/release`; machine-local) | shares local + remote box; flock-serialized; exclusivity-checked | When the task fits one of these families, EXTEND it (add a shape row / a mode), don't write a new harness. Only author a fresh `test_*_faithful_pertask.py` for a genuinely new family (e.g. attention) — and even then mirror the group-GEMM file's structure exactly. ## Recipe — stand up + run a faithful gate for a new (kernel, shape) 1. **Find the task_type_id.** Get the numeric `TASK_*_SM100` enum from `runtime_header.h` / the kernel's `register_*` in `src/kernel/task_register.cc`. Match the profiler row by this NUMERIC id, NOT the name — many task names are missing from `profiler_persistent.py`'s `event_name_list` so the CSV writes `UNKNOWN_<id>`. `profiled_per_task_latency_by_id(pk, TASK_ID, iters, label=...)` handles this. 2. **Build the input at PRODUCTION decode geometry** (see "Decode geometry" below) — this is the #1 correctness trap. Register the kernel through `PersistentKernel(test_mode=True, num_workers=resolve_num_workers(...))`, attach inputs, `pk.compile()`, `pk()` once for correctness, then the profiled timing run. 3. **Report** slowCTA + wall + sumCTA + cos, and emit a single `FAITHFUL_RESULT {...}` JSON line (mirror the group-GEMM gate's `_lat_table` / FAITHFUL_RESULT format) so `faithful_eval.py` can parse baseline vs candidate. 4. **Per-worker-count sweep.** The optimal kernel FLIPS with the grid (GEMV wins at high occupancy nw≥128; tcgen05 MMA wins at low nw≤68). Sweep **nw ∈ {8, 64, 68, 128, 136}** (128 first), via `MPK_TEST_NUM_WORKERS` / the `--num-workers` plumb. The 64/68 points feed the group-GEMM ‖ shared-expert worker-partition design. 5. **Run it via the broker** (never pick a fixed GPU; faithful needs an EXCLUSIVE card): ``` ~/gpu_broker/gpu_gate.sh --holder <tag> -- \ --kind <finen|gemv_m1|largem_compact> --shape <S> --num-workers <W> \ --configs <S> [--candidate-role smallm] [--kernel <cand.cuh>] [--baseline] [--decode-m1] ``` It auto-acquires a free local-or-remote slot, runs the gate, streams `KERNEL_RESULT`/`FAITHFUL_RESULT` back, and releases (trap-on-EXIT). `--baseline` = the in-tree kernel (ratio 1.0); `--kernel <cand.cuh>` overlays a candidate via a shadow MIRAGE_ROOT (production tree never written). ## Decode geometry (the critical correctness point) Production decode is **bs=1 → M=1 per active unit**. A gate that feeds an all-128-rows-real activation measures a PREFILL-like geometry and CANNOT reward an M=1-aware kernel. - **Dense projections**: M=1 (the dense gate is already M=1). - **Routed group-GEMM**: exactly 1 real token per ACTIVE expert (top-k routing). Build `a_bf16` with only the live row (`expert*128 + 0`) of each active expert non-zero (the other 127 rows are pad/zero), and validate cos on the LIVE rows ONLY. This is the `MPK_GG_DECODE_M1=1` mode. The active experts come from `meta` row-1's `active_expert_mask` (~4–8 active at bs=1 TP8 EP2). - **MLA / attention**: per-rank decode (1 query token, the real KV-seq length). ## Invariants & gotchas (these have each cost real hours — bake them in) 1. **Exclusive GPU only.** A faithful slowCTA on a contended card is INVALID. The gate torch-probes + exclusivity-checks (refuses on a foreign compute proc). The broker free-checks before acquiring. 2. **Gate watchdog (`MPK_GATE_TIMEOUT_S`, default 420s).** A buggy candidate can DEADLOCK the megakernel and wedge the (exclusive) box indefinitely. `faithful_eval.py`'s `_run_test` wraps the test in `subprocess.run(timeout=...)` → on timeout it SIGKILLs the child = the megakernel-HOST process → frees the card → returns rc=124 `VERDICT=TIMEOUT_HANG`. The timeout MUST live in the .py (which runs the test as its subprocess); a wrapper-level `timeout` orphans the test grandchild (it keeps holding the GPU). Keep this. 3. **Candidate kernels: NO block barrier inside the persistent-sweep tile loop.** In MPK each worker strides tiles by `num_workers`; workers with fewer tiles EARLY-EXIT the loop, so a `__syncthreads()`/`bar.sync` inside it deadlocks (the exited workers never reach it). Each lane/warp does its full K-sweep independently; only a final intra-warp `__shfl` reduction is allowed. (This deadlock wedged the exclusive box for 24 min.) 4. **Remote box runs its OWN repo copy.** The exclusive remote box (`<BOX_USER>@<BOX_IP>` — site-specific, resolve per session per `v2-model-support/references/box-orchestration.md` §1-2; its mirage repo + ferret-scripts paths are box-local) does NOT see your local edits. After editing a gate test / `_build_helper.py` (new shape) / `faithful_eval.py` (watchdog, new flag), `scp` them there (keep a `.bak`). Verify with a remote `python3 -c "import ast; ..."`. 5. **The remote box's `faithful_eval.sh` has an arg-allowlist** (and is OFF-LIMITS to edit — along with `cc-run.sh`). It REJECTS unknown flags (`--decode-m1` → "unknown arg"). Forward a new knob as an ENV var instead, via the machine-local box shim (`~/nebius_gate/faithful_eval_<box>.sh`, which translates e.g. `--decode-m1` → `MPK_GG_DECODE_M1=1` prefix on the remote command). The test reads the env var directly and `_run_test` does NOT pop it. Local-only: the ferret `faithful_eval.sh` similarly hardcodes a stale "GPU 5 = KDA job" refusal — route around it via the broker `gpu_pool.conf` (disable the blocked local slot), don't edit the .sh. 6. **Add `--num-workers` / new flags to `faithful_eval.py`** (argparse → `run()` → `_run_test`/`_run_group_gemm` → the test env), NOT just an env var, so they cross the remote shim as args. (The shim then env-translates the ones the remote .sh rejects.) 7. **Per-(shape) one number.** The gate measures ONE shape per invocation; mapping it onto multiple differently-shaped `--configs` mis-gates. Pass `--configs` matching `--shape`. ## Dispatching a kernel agent against the gate When handing a kernel to KDA/Ferret (the `kda-kernel-agent` / `ferret-kernel-agent`), the faithful gate IS the promotion authority. Give the agent: the exact `gpu_gate.sh` command (with `--kind/--shape/--num-workers/--decode-m1`), the measured in-tree baseline slowCTA (measure it FIRST), the vLLM ref + the ≥20% target (`≤ vLLM_ref ÷ 1.2`), the per-worker-count sweep set, and the barrier-free-tile-loop + alignment crash-safety constraints (gotcha #3). The agent's candidate is overlaid via `--kernel`; the shadow MIRAGE_ROOT keeps the prod tree clean. ## Output: the FAITHFUL_RESULT contract Each gate run prints one machine-parseable line the bridge/agents consume, e.g.: ``` FAITHFUL_RESULT {"kind":"...","shape":"...","num_active":8,"candidate":{"slowCTA_us":..,"wall_us":..,"sumCTA_us":..,"cos":1.0,"status":"PASS"},"baseline_in_tree":{...},"slowCTA_ratio_base_over_cand":..,"verdict":"PASS_GATE|BELOW_BAR|TIMEOUT_HANG|INVALID"} ``` Keep this shape stable — the ferret/KDA loop greps it.