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