AKO4X

Run performance benchmarks to get a verdict on whether a kernel change actually helped. Use --ab-compare for sub-1x deltas (drift cancels in same container), --variance-check for noise floor, and the subset modes (--first, --smoke; --extremes is modal-backend only) for compile-correctness probes only — those are NOT performance verdicts. The only authoritative source for "is this change real". Benchmark specifics (config.toml schema, status enum, scoring formula, baseline rule, fresh-inputs contract) live in the `benchmark` skill.

Reference for the active benchmark harness — what it IS and how it behaves (the active benchmark is flashinfer-bench). Covers config.toml structure (`[solution]`/`[build]`/`[benchmark]` tables), the status enum, the workload / fresh-inputs model, reference-baseline + scoring mechanics, tolerance keys, TVM-FFI builder link limits, and the silent-skip cascade. Invoke whenever you decode a bench status string, hit an unfamiliar config.toml field, suspect "correctness passed but the headline is implausible", or need to know what's frozen before proposing a bench edit — do NOT guess field or status semantics. Bench commands and noise methodology (A/B compare, variance check, drift cancellation) live in the `bench` skill.

Host-side C++ kernels (`[build].language = cpp`) — the `solution/{kernel.cpp, binding.py}` two-file TVM-FFI binding pattern, and when to reach for it (host-side orchestration of a kernel you write, or composing CUTLASS headers into it). Distinct from the `cuda` skill which covers `.cu` device kernels compiled by nvcc. Use when writing a host-side C++ kernel or deciding cpp-vs-cuda.

CUDA C++ (.cu) kernel reference — TVM-FFI direct-export vs Python-binding entry points, the chevron-launch null-stream pitfall under CUDA-graph capture, the sm_100 fp8→bf16 cvt PTX gap, the load_inline name-cache trap, and the __launch_bounds__ register-spill lever. Also the single-source for generic per-call-overhead / CUDA-graph-capture / PDL theory (Waves-Per-SM decision table) regardless of DSL. Use when writing or debugging hand-written CUDA C++ kernels, or whenever you need generic CUDA-graph / sync-audit / PDL reasoning.

CuTe DSL reference — NVIDIA's CUTLASS Python DSL. Use whenever writing or debugging a CuTe DSL kernel — covers the `@cute.kernel` + `@cute.jit` + `.launch()` host pattern, `from_dlpack` tensor conversion, API-probing against the installed wheel, the launch-kwarg table, and the well-known graph-capture pitfall (CuTe's TVM-FFI stream binding does NOT pick up CUDA-graph capture mode — flashinfer-bench issue 414).

Run NVIDIA Nsight Compute (NCU) per-kernel profiling for register pressure, occupancy, stall-reason, memory-throughput, and IPC analysis. Use BEFORE architecting an optimization fix — not only after — when a hypothesis about microarchitectural behavior needs verification. NCU `Duration` is NOT comparable to bench timing; use NCU for ratios only.

Run NVIDIA compute-sanitizer (memcheck / racecheck / initcheck / synccheck). FIRST action on `INCORRECT_NUMERICAL` or flaky output — rolling back destroys the evidence; sanitizer often points at the exact line. Also useful for diagnosing race conditions in multi-block kernels and uninitialized-memory reads.

TileLang DSL reference — the `@tilelang.jit` factory + `@T.prim_func` pattern, and the TileLang PDL binding (`T.pdl_trigger()` / `T.pdl_sync()`, with JIT auto-setting the launch attribute so no host-side flag is needed, unlike Triton). Use when writing or tuning a TileLang kernel, or wiring TileLang↔Triton PDL overlap.

Triton DSL reference for kernel writers — num_warps/num_stages choice, the small-N fp8 MMA-throughput regression, split-K reduce tile form, deterministic tl.join/permute concat (tl.cat is not order-stable), autotune + .triton_cache pitfalls, and the Triton PDL binding. Use whenever writing or tuning a Triton kernel; first stop before guessing tile shapes or pipeline depth.