mirage-project
- 24 skills
- 0 followers
- 6 hours ago last updated
- ▌ Test Mode · mirage-projectGuide for using MPK test mode to unit-test individual layers or multi-layer pipelines through the full compilation pipeline. Use when writing layer tests, debugging kernel output, or validating a new task end-to-end.
- ▌ Add Mpk Task · mirage-projectStep-by-step guide for adding a new task implementation to Mirage Persistent Kernel (MPK). Use this when adding a new GPU operator (e.g., a new attention variant, normalization, activation) to the MPK megakernel.
- ▌ Add Mpk Model · mirage-projectGuide for adding a new model (e.g., Llama4, DeepSeek V3) to the MPK persistent kernel. Covers prerequisites check, demo structure, layer wiring, and testing.
- ▌ Mpk Internals · mirage-projectReference guide for the MPK compilation-to-runtime pipeline. Use when asked how MPK works internally, how compilation/code generation works, what happens at runtime, or when debugging the megakernel scheduler.
- ▌ V2 Model Support · mirage-project bundleEnd-to-end pipeline for adding or porting a model to MPK Runtime-V2 — from a compute-graph spec (shapes + draw.io graph + HF checkpoint + TP/EP plan) to a working multi-GPU demo. Use when bringing up a NEW model on the v2 (role-split, static-plan) runtime, when porting an existing v1 model to --use-v2, or when handed a compute-graph file and asked to make it run. Covers graph→plan, builder/demo bring-up, per-kernel authoring dispatch, the debug gate ladder, and the multi-agent/box workflow.
- ▌ Mpk Faithful Gate · mirage-projectBuild or run a FAITHFUL in-MPK per-task latency gate (slowCTA at the production grid + cos) for a DeepSeek-V3 MPK decode kernel or shape. Use this WHENEVER you need to measure, gate, or head-to-head-optimize an MPK kernel's per-task latency (dense FP8 GEMM, routed group-GEMM W13/W2, MLA decode, router/topk, AllReduce, etc.), stand up a faithful gate for a NEW kernel/shape, or dispatch a KDA/Ferret kernel agent against a faithful measure — and ESPECIALLY before trusting any per-task µs number in a DSv3 decode perf campaign. The faithful in-MPK slowCTA is the trusted measure; a standalone green-ctx bench MIS-RANKS and a whole-megakernel e2e number hides per-task cost — do not use either as the gate. Covers the slowCTA definition, the _faithful_helper reuse, the decode (M=1) input geometry, the GPU-broker + exclusive remote box, the gate watchdog, and the candidate-overlay bridge to faithful_eval.py.
- ▌ Mpk Lever Cleanup · mirage-projectUse when a batch of env-gated (`#ifdef MPK_DSV3_*` / `os.environ`-controlled, default-OFF) MPK optimization levers needs to be consolidated into a single clean code path for a PR: hard-wire every winning lever as the default, delete the legacy `#else` branches, remove the env vars that select new-vs-old logic, revert dead levers that measured KILL/NULL/regress, delete diagnostic probes, then commit one clean version. Applies to the wrap-up stage where the optimization work has settled and is being merged to mainline. Not for the exploration phase (levers should stay env-gated default-OFF there) or for runtime/execution-model changes.
- ▌ V2 Kernel Writing · mirage-project bundleRuntime-V2 kernel-writing workflow. Use when writing, porting, or rewriting ANY Runtime-V2 task kernel (tasks/blackwell_v2/*.cuh + registration) — a new op, a v1→v2 port, or a rewrite toward the reference linear_sm100_v2 warp-role pipeline idiom. Drives the staged loop SPEC→IMPLEMENT→WIRE→VALIDATE→PERF→REVIEW with per-stage subagents and the b200-* sub-skills, and enforces the M=1 anti-loop evidence + the v2 protocol invariants (§1.1 dep-prefix, stale-arrival re-init, skip_after_step0, task_offset wiring).
- ▌ V2 Perf Iteration · mirage-project bundleRuntime-V2 performance-iteration workflow. Use when running a perf-optimization campaign or iteration on the v2 runtime (--use-v2) — measuring a baseline, ranking bottlenecks, planning levers, implementing, re-measuring, and landing/recording the verdict. Drives the loop MEASURE→ANALYZE→PLAN→NEXT-MOVE→REVIEW→IMPLEMENT→VALIDATE+RE-MEASURE→LOOP-OR-LAND→RECORD with the mpk-* subagent roster, the v2 profiler/perfetto toolchain, and the TIER-1 TP8 verdict discipline.
- ▌
- ▌ Mpk Development Norms · mirage-project bundleThe MPK team's "where does a change belong + what a clean PR looks like" norms, extracted from mirage-project/mpk merged-PR history. Read FIRST — before starting any MPK change, opening/shaping a PR, deciding which file a change goes in, reviewing a diff's shape, or cleaning up a branch that grew messy/off-norm ("改得太乱/不符合开发规范"). Complements add-mpk-model / add-mpk-task / v2-model-support (the HOW) with the WHERE + the PR-shape gate.
- ▌ Dpskv3 Logistic Review · mirage-projectAudit the DeepSeek V3 MPK demo + builder chain end-to-end and confirm logical equivalence with vLLM's reference implementation. Use after structural changes to `python/mirage/mpk/models/deepseek_v3/builder.py`, `demo/deepseek_v3/demo.py`, or any MLA / MoE / MTP task in `python/mirage/mpk/persistent_kernel.py`, `src/kernel/task_register.cc`, and `include/mirage/persistent_kernel/tasks/blackwell/mla_*.cuh` / `moe_*.cuh`. Produces a structured drift report so the change reviewer can confirm the math/topology is still equivalent.
- ▌ B200 Scope Layout Dispatch · mirage-project bundleUse when the user wants to map an ML operator onto a B200/Blackwell kernel, or to review which threads should execute a given tile primitive, where the data should live, and whether to invoke thread code or TMA/tcgen05. Outputs a complete contract for scope, layout, dispatch, and handoff. Not for high-level model architecture design alone or pure hardware spec queries.
- ▌ B200 Tma Pipeline Designer · mirage-project bundleUse when the user wants to convert regular GMEM↔SMEM tile copies to TMA on B200/Blackwell, design double-buffered/multi-stage pipelines, choose a swizzle, or distinguish the TMA load vs store completion protocols. Produces descriptor, stage ring, barrier, and prologue/steady-state/epilogue plans. Not for highly irregular gather/scatter or copies too small to be worth setting up a TMA descriptor.
- ▌ B200 Kernel Roofline Triage · mirage-project bundleUse when the user asks "why is this B200/Blackwell kernel slow, what should I optimize first, should I fuse or move to Tensor Cores/TMA". Classifies the problem via arithmetic intensity, dataflow, and resource occupancy as memory-bandwidth-, compute-throughput-, latency/concurrency-, or scheduling-bound, and gives a minimal falsifiable experiment. Not for queries that only ask about hardware specs with no kernel/operator context.
- ▌ B200 Tmem Lifecycle Planner · mirage-project bundleUse when the user is using `tcgen05` on Blackwell/B200 and needs to plan TMEM accumulators, block-scale factors, TMEM column allocation, the `tcgen05.ld/st/cp` paths, epilogue readback, and safe deallocation. Produces a TMEM region map, column budget, and lifecycle. Not for Hopper/Ampere or code that treats TMEM as ordinary shared memory.
- ▌ B200 Layout Contract Auditor · mirage-project bundleUse when a B200/Blackwell kernel shows wrong results, uncoalesced global memory access, SMEM bank conflicts, a TMA swizzle that mismatches the Tensor Core read, or confused TMEM/register ownership. Audits shape–stride, thread distribution, swizzle, and the hardware operand contract layer by layer. Not for pure synchronization deadlocks or compile errors unrelated to memory layout.
- ▌ B200 Flash Attention4 Planner · mirage-project bundleUse when the user wants to design or extend a FlashAttention-style forward kernel on B200/Blackwell, involving the two MMAs QKᵀ and PV, online softmax, S/P/O in TMEM, warp roles, causal mask, GQA, tile scheduling, or final normalization. Outputs the algorithm state, tile graph, barrier graph, and validation plan. Not for cases that only use off-the-shelf framework operators, where the full backward is not yet defined, or for ordinary dense GEMM.
- ▌ B200 Gemm Optimization Ladder · mirage-project bundleUse when the user wants to implement from scratch, port, or systematically optimize a B200/Blackwell GEMM. Advances level by level along "correct single tile→K loop→spatial tiling→TMA→multi-stage pipeline→persistent→warp specialization→2-CTA cluster→multi-consumer", with a correctness and performance gate at every level. Not for cases that only want to call a mature BLAS and need no custom fusion/layout.
- ▌ B200 Mbarrier Protocol Auditor · mirage-project bundleUse when a Blackwell/B200 asynchronous kernel deadlocks, fails intermittently, reads stale data, reuses a stage too early, or when the arrival, tx-count, phase, and wait of TMA/tcgen05/CLC need auditing. Produces a per-barrier protocol ledger and fix points. Not for cases already known to be out-of-bounds accesses, layout mismatches, or ordinary host-side synchronization problems.
- ▌ B200 Warp Specialized Debugger · mirage-project bundleUse when a B200/TIRx/CUDA warp-specialized kernel fails to compile, deadlocks, hits an illegal memory access, produces wrong results, or is "correct but slow". First verifies the environment and a minimal reproduction, then builds a roles/storage/handoff/lifetime worksheet from the generated CUDA/PTX, fixing one handoff at a time. Not for cases with no reproducible code yet or purely high-level model accuracy issues.
- ▌ B200 Cluster Persistent Scheduler · mirage-project bundleUse when the user wants to use Thread Block Cluster, DSMEM, 2-CTA cooperative MMA, persistent kernels, a tile scheduler, or Cluster Launch Control on B200 to handle uneven tails. Outputs the cluster tile, occupancy, sharing/multicast, and static or dynamic scheduling plan. Not for kernels where independent small CTAs already saturate the machine and the workload is uniform.
- ▌ B200 Tcgen05 Mma Contract Builder · mirage-project bundleUse when the user needs to choose the tile, dtype, `cta_group::1/2`, SMEM operand layout, or TMEM accumulator mapping for a `tcgen05` MMA on B200/Blackwell, or to implement an mxfp8/nvfp4 block-scaled GEMM. Produces an auditable MMA contract and completion protocol. Not for ordinary CUDA-core matmul or non-Blackwell targets.
- ▌ Blackwell Build Compatibility Auditor · mirage-project bundleUse when the user wants to confirm whether an existing CUDA extension/binary can run on B200, configure `compute_100/sm_100` or the architecture-specific `sm_100a`, or check PTX/cubin, CUDA Toolkit versions, JIT, and fatbin. Outputs compatibility evidence, build flags, and a smoke test. Not for kernel performance tuning once the kernel already runs correctly.