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.