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)

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.