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