V2 Kernel Writing

You are writing a Runtime-V2 task kernel for MPK (dsv3-decode-clean). The house style is the reference linear_sm100_v2.cuh warp-role pipeline (loader W4 TMA → launcher W5 tcgen05/TMEM → consumers W0-3 epilogue → storer W6 page release), with consumer-only as the sanctioned idiom for non-GEMM-shaped ops. Quality bar and every protocol invariant live in this skill's references — the loop below tells you when to read what and what to dispatch.

For MODEL-level bring-up (whole compute graph → demo) use v2-model-support; this skill is the per-KERNEL inner loop that pipeline dispatches into.

Environment prerequisites (what must exist on the machine)

Everything the loop REQUIRES travels with the repo: this skill's references/ + applications/ docs, the reference kernels (include/mirage/persistent_kernel/tasks/blackwell_v2/), the harness (tests/runtime_python/blackwell_v2/), and the wiring surface. The rest is machine-local and DEGRADES GRACEFULLY:

• b200-* sub-skills (the Skill-tool names in the index below) are USER-global (~/.claude/skills/b200-*), NOT in this repo — a fresh clone still has them only under the same user account. If missing: proceed anyway; references/house-style.md + references/upstream-kernel-catalog.md carry the distilled protocol/layout contracts.
• Upstream catalog reads (git show mirage-project/runtime_refactor:<path>) need the remote: git remote add mirage-project https://github.com/mirage-project/mirage.git && git fetch mirage-project runtime_refactor. Optional — the catalog doc is self-contained.
• ~/ferret/ (TEMPLATE_v2.yaml, docs/v2_runtime_notes.md, workspace1..8) — required ONLY for the ferret-v2 engine; ~/kda-workspaces/ only for kda; ~/kernel_tools/ (ncu_profile.sh) only for the NCU bound-check. Absent ⇒ route Stage 2/5 work to v2-kernel-engineer (the default anyway) and use b200-kernel-roofline-triage/manual NCU.
• Personal memory (~/.claude/projects/-home-muhengl-mirage/memory/) — optional context only (same-user machines). references/m1-decode-evidence.md is the self-contained distillation of the evidence rows; treat the memory dir as its (optional) source citations.
• Hardware: Stages 0-3 need no GPU. Stage 4/5 need a B200 (sm_100a); TP8-geometry ops need the 8-GPU box (see v2-model-support/references/box-orchestration.md) — without it, deliver with local gates + an explicit "pending TIER-1".
• Codex MCP (mcp__codex__codex) for Stage-6 double-checks — if unconfigured, the ablation-logic-reviewer pass still runs; note the missing second engine in the verdict.

Reference docs (this skill's folder)

Dispatch pattern (subagent nesting)

The MAIN THREAD (or one lead subagent for a multi-kernel campaign) is the orchestrator: it runs Stages 0/3/6 itself and dispatches ONE subagent per heavy stage — a designer (Stage 1), an implementer (Stage 2 — use the v2-kernel-engineer agent if defined in .claude/agents/, else general-purpose with that discipline pasted in), a validator (Stage 4). Subagents do not dispatch subagents. Each dispatch prompt MUST name: the stage's reference docs (absolute paths), the sub-skills to load via the Skill tool, the op contract, and the exact deliverable. Stages are sequential; iterate 2↔4 on failures, 5→1 on a perf verdict that changes the design.

Hard rules for every stage: default build byte-identical (new task types additive; levers env-gated default-OFF); no repo-wide refactors; GPU safety (test-mode first, never crash-loop the megakernel); every non-trivial conclusion → Stage 6 review before acting on it.

Stage 0 — LOAD (orchestrator, no code)

Read references/house-style.md + references/m1-decode-evidence.md IN FULL before any design. Then classify the op: shape (M, N, K / attention / elementwise), dtype, per-token work, TP/EP sharding, where it sits in the layer DAG (producer/consumer events), and which evidence rows (D*/W*/U*) touch it. Match the op to an upstream family (A pipeline / B consumer+regions / C consumer+monolith / D consumer+no-SMEM / E sub-op helper) in references/upstream-kernel-catalog.md — the family names the file to crib from. If the op matches a DEAD row and no new mechanism is on offer — STOP and say so; that is a successful outcome of this stage.

Stage 1 — SPEC (designer subagent)

Deliverable: a spec.h-style design doc (markdown). Location convention: campaign items that should travel with the repo go to .claude/skills/v2-kernel-writing/applications/<item>_spec.md (exemplar: applications/ffn_item1_spec.md); throwaway/exploratory specs go to scratch/ (git-ignored, machine-local). The spec contains:

1. Engine choice via the DECISION TREE:
   • GEMM-shaped & M ≥ tile (or a real K-pipeline: ≥ 4 K-iterations of streamed weight tiles; page count per tile is dtype-dependent — an fp8 128×128 tile is exactly 1 page and still qualifies) → reference pipeline w/ TMA + tcgen05. Load b200-tma-pipeline-designer (stage ring, swizzle, load-vs-store completion), b200-tcgen05-mma-contract-builder (tile/dtype/ cta_group, SMEM operand layout, I-desc), b200-tmem-lifecycle-planner (TMEM columns, alloc/dealloc, ld/wait). Anchor every choice to house-style §2/§5.
   • M=1 GEMV / memory-bound streaming → consumer-GEMV per m1-decode-evidence (D2/D3/D9); nwarps per the W1 wave-quant tie-test ceil(items/(ntasks*nwarps)).
   • Attention-shaped → consumer-only per upstream attention_sm100.cuh (house-style §0); for FA-style rewrites additionally load b200-flash-attention4-planner.
   • Consumer-only ops also pick a SMEM family (catalog §2-4): planner-region typed buffer struct (rmsnorm/argmax pattern — the DEFAULT for anything staging through SMEM); ONE monolithic honest region when porting a kernel with hand-rolled internal offsets (attention pattern — never declare regions the device won't address); NUM_REGIONS=0 spec for pure-GMEM elementwise (silu/embedding pattern). A tiny fused sub-op (norm/rope-style) may be a family-E SMEM-view helper inside a host kernel instead of a new task.
   • Unsure how to map the op at all → load b200-scope-layout-dispatch first.
2. SMEM region plan + budget: named regions, sizes, can_pack, page math (16KB pages, ≤14 pages, total ≤ 224256 B), alignment 1024 — house-style §4 format.
3. SEM ordinal table (house-style §3 format: ordinal, count, producer→consumer, meaning; ≤31 op-private) — or the tag-flag alternative (W4) with its flag layout, if multi-role handshakes on the consumer-only idiom.
4. Stage count + role responsibilities per role, incl. who re-inits which async mbars and who releases which pages on which path (bounds-fail included).
5. Task granularity: DEFAULT per-tile (tile_idx = task_offset). Grid-wide fused is the EXCEPTION — requires written justification + the monotonic-barrier + skip_after_step0 + num_tasks==num_workers contract (house-style §6, wiring-recipe §7/§8).
6. Evidence check: per design choice, the D*/W*/U* row it rests on; for U* probes, the pre-registered predicted Δ + kill threshold.

Stage 2 — IMPLEMENT (implementer subagent)

Engine choice first (routing rule + flow in references/ferret-v2-dispatch.md): v2-kernel-engineer = protocol-heavy house-style port / first bring-up (default); ferret-v2 (ferret-kernel-agent V2 MODE) = beat-a-numeric-TIER-2-target optimization loop once a faithful FROZEN gate exists (op wired + harness case + live anchor) — also the "ferret writes from spec" path over a wired stub; kda = verdict-grade honest transfer when the number decides a campaign verdict and over-claim is costly. Ferret-v2 rearranges the stages (S3-stub wiring precedes the run — the gate substrate is the in-tree harness); its friction escape falls back to engineer-shell + ferret-math-only-body.

Write <op>_v2.cuh + <op>_v2_spec.h to the spec. House-style code conventions:

• namespace kernel { namespace <op>_v2 {; spec.h constants + static_asserts pinning every mirrored constant; role-split __device__ __noinline__ functions (one per role).
• mbarrier protocol per house-style §2/§3: start-of-task re-init of async-arrived mbars by their arriving role; fence.mbarrier_init.release.cluster after inits; tcgen05.fence::after_thread_sync at MMA↔wait boundaries. Before finalizing, invoke b200-mbarrier-protocol-auditor on the barrier ledger (every mbar: init count, arrivers, tx-bytes, waiters, phase evolution, re-init ownership).
• extern __shared__ __align__(1024); SMEM only via task_desc->smem_region_offset(REGION_*).
• No __syncthreads() in role loops — named barriers (bar.sync <free-id>, 128) or tag-flags only; elect_sync() for single-thread issue; no blockIdx for identity.
• Layout doubts (swizzle vs tcgen05 operand, coalescing, bank conflicts) → b200-layout-contract-auditor. Build-flag doubts (sm_100a, -rdc=true) → blackwell-build-compatibility-auditor.

Stage 3 — WIRE (orchestrator)

Follow references/wiring-recipe.md top to bottom — enum, task_header include, register fn (§1.1 dep-prefix is the first line of the consumer body — MANDATORY), graph.cc tuple, runtime.cc task_type_to_name + task_offset=bid.x list, py wrapper (num_tasks==num_workers gate if grid-wide), builder use_v2 branch, skip_after_step0 on any monotonic-barrier scratch. Tick the §10 ship checklist explicitly.

Stage 4 — VALIDATE (validator subagent + references/validation-debug.md)

In order, no skipping: (1) test-mode numeric vs torch in tests/runtime_python/blackwell_v2/ (cos ≥ 0.999, rel_max ≤ 3e-2, no NaN; v1-counterpart compare); (2) §1.1/protocol static audit; (3) in-MPK --layers 0-3 probe; (4) MULTI-STEP run, iter ≥ 3 — iter-0-fine/iter-1-hang = persistent-state re-init (skip_after_step0), NOT a missing event; (5) on any hang: watchdog (-DMPK_V2_BREADCRUMB + MPK_V2_HANG_WATCHDOG_S); on any crash: compute-sanitizer memcheck = ground truth (breadcrumb counts are base-rate-biased). Deadlock/wrong-result debugging → b200-warp-specialized-debugger (roles/storage/handoff/lifetime worksheet, one handoff at a time). Math-changing on TP8 → poison-fill gate, not token-identity.

Stage 5 — PERF (orchestrator or validator)

TIER hierarchy is the law: TIER 1 in-MPK %globaltimer slowCTA @ production grid = the only verdict-grade number; harness slowCTA corroborates; cudaEvent-wall / standalone-warm are diagnostic-only. Compare against the reference/v1 body anchor from the spec. Bottleneck classification → b200-kernel-roofline-triage (achievable-floor rules from m1-decode-evidence §4 apply — same-grid xor-consumer floor, never theoretical peak). For a pipeline kernel that is correct-but-slow, climb b200-gemm-optimization-ladder one rung at a time. Profiler: buffer = 120000*128 entries; export via scripts/v2_perfetto_export.py. For a sustained beat-a-numeric-target optimization loop on one kernel, dispatch ferret-v2 (references/ferret-v2-dispatch.md): it iterates the pair against the frozen harness gate (TIER-2 body_span) autonomously; TIER-1 in-MPK slowCTA stays the final verdict here. For a whole perf-optimization CAMPAIGN around this kernel (measure→plan→implement→re-measure →land, agent roster + history contract) use the sibling skill v2-perf-iteration.

Stage 6 — REVIEW (orchestrator)

• EVERY non-trivial conclusion (root-cause, DEAD/ALIVE verdict, perf claim, "matches reference") → ablation-logic-reviewer subagent + Codex MCP double-check (default params) BEFORE acting on or reporting it.
• Landing: mpk-correctness-gate for anything math-adjacent, then mpk-commit-reviewer before git commit (staged-path + byte-identity + message gates). Verdicts → mpk-memory-keeper (experiment_history INDEX + memory; update m1-decode-evidence sources).

Sub-skill index (load via Skill tool, exact names)