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).