mirage-project/mirage

add-mpk-model

Guide for adding a new model (e.g., Llama4, DeepSeek V3) to the MPK persistent kernel. Covers prerequisites check, demo structure, layer wiring, and testing.

add-mpk-task

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

b200-cluster-persistent-scheduler

Use 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…

b200-flash-attention4-planner

Use 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…

b200-gemm-optimization-ladder

Use 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…

b200-kernel-roofline-triage

Use 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-,…

b200-layout-contract-auditor

Use 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…

b200-mbarrier-protocol-auditor

Use 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.…

b200-scope-layout-dispatch

Use 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,…

b200-tcgen05-mma-contract-builder

Use 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…

b200-tma-pipeline-designer

Use 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…

b200-tmem-lifecycle-planner

Use 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…

b200-warp-specialized-debugger

Use 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…

blackwell-build-compatibility-auditor

Use 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,…

dpskv3-logistic-review

Audit 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 /…

ferret-kernel-system

Use when you need a NEW or optimized MPK CUDA kernel (a per-task `.cuh` under include/mirage/persistent_kernel/tasks/) that must PROVABLY beat a target WITHOUT the kernel-optimizer cheating its own correctness test. This is the frozen-gate Ferret system — a…

mpk-development-norms

The 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…

mpk-faithful-gate

Build 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,…

mpk-lever-cleanup

Use 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`…

v2-kernel-writing

Runtime-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…

v2-model-support

End-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,…

v2-perf-iteration

Runtime-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.…

test-mode

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

mpk-internals

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