What problem does it solve? DeepSeek-class MLA decode attention on AMD MI300X (gfx942) is memory-bound on the KV-latent read, and the live TileLang kernel is hard to autotune and carries host-sync / CUDA-graph-capture hazards that block further optimization. ## Core Features & Use Cases - TileLang-to-Triton porting procedure: Step-by-step recipe to reimplement the MLA decode core in Triton with flash-style online softmax and absorbed up-projections (W_UK/W_UV) so the latent never expands to full head_dim in HBM. - Graph-capture safety rules: Enforces a host-sync-free hot path (no .item()/.cpu()/.synchronize()) with weight prep cached by data_ptr, so the kernel can overlay the live decode path. - Tuning and parity guidance: Concrete knob ranges (BLOCK_M=16/32, split_k, num_warps, fp8 fused scaling) plus do-no-harm pitfalls such as split_k numerical errors and bf16 weight re-materialization memory blow-up. - Use Case: While running the GEAK kernel_workflow on a DeepSeek model's decode bottleneck, apply this advisory prior to port the TileLang MLA kernel to a tunable, capture-safe Triton version targeting at least 1.15x isolated speedup with required parity. ## Quick Start Ask the agent to apply the mla_tilelang_to_triton expert skill to port the TileLang MLA decode kernel to a graph-capture-safe Triton kernel on gfx942 and validate parity against the TileLang oracle.
