tilegym-improve-cutile-kernel-perf

Create a git branch, locate the kernel under src/tilegym/suites or src/tilegym/ops/cutile, classify it as memory-bound or compute-bound, then run the experiment loop: apply one optimization per iteration, verify correctness with pytest, benchmark latency, and record results in perf_results.md.

Replacing gather/scatter with TMA loads and stores yields 2-78x gains, persistent scheduling adds 50-300% for many work items, and autotuning tile sizes, occupancy, and num_ctas adds 10-50%. The optimization playbook lists recipes A through L in priority order.

Benchmarking requires a GPU node with Blackwell, Hopper, or Ampere architecture. Note that Ampere lacks hardware TMA, so ct.load with allow_tma falls back to cp.async emulation and needs padding_mode=ZERO to avoid silent out-of-bounds corruption.

Common causes are flush_to_zero or APPROX rounding changing results, tile size out-of-bounds access, allow_tma=False semantics, and persistent loop bound errors. The skill reverts any change that fails correctness immediately and moves to the next optimization.

Use IR and SASS analysis when performance is unexpectedly poor but results are correct, to inspect scheduling, MUFU instruction counts, and token dependency chains. Do not use it for compile errors, wrong numerics, or differences under 5%.