GPU Development Skills

---
name: tilelang-optimization-loop
description: Run a bounded TileLang implementation and optimization task from a PyTorch reference or existing kernel through baseline analysis, a written plan, candidate edits, lowering inspection, validation, measurement, and promotion. Use for iterative kernel work and reproducible handoffs; use a specialist directly for a single API, debugging, or profiling question.
---

# TileLang Optimization Loop

Follow a repeatable experiment cycle in the **task workspace**. Keep the skill
suite and its upstream knowledge separate from generated kernels and results.
Resume from evidence that still matches the source, inputs, and environment.
This adapts the [KDA workflow](https://github.com/mit-han-lab/kernel-design-agents/blob/4806866492d5c7cad05e07918790bd1e94e976a1/docs/agent-flow.md)
to the chain from TileLang source through compiler lowering to GPU behavior.

## Route by the current decision

At each stage, name the question, read the selected specialist's entrypoint,
and consume its finding in the next decision. A link alone is not a completed
specialist step. Reuse matching evidence rather than rereading or rerunning
everything. Use the [stage report](references/stage-handoff.md) for iterative
tasks and session/worker handoffs; link it from the existing candidate ledger.

| Current question | Skill to read |
| --- | --- |
| What does the PyTorch baseline execute and allocate? | [Torch profiling](../tilelang/torch-profiling-tilelang-programs/SKILL.md) |
| Which hardware mechanism or instruction contract informs this choice? | [KernelWiki interface](references/kernelwiki.md), then the installed KernelWiki skill |
| Does a high-level placement or fusion choice need a semantic/static model? | [TileFoundry design](../designing-with-tilefoundry/SKILL.md) |
| How should this candidate express the operation? | [Writing kernels](../tilelang/writing-tilelang-kernels/SKILL.md) |
| Why is output wrong, or why did the intended lowering disappear? | [Debugging](../tilelang/debugging-tilelang-programs/SKILL.md) |
| Are the gradients and saved-state contract correct? | [Forward/backward testing](../tilelang/testing-fwd-bwd-kernels/SKILL.md) |
| Is the comparison reliable, and what does the profile show? | [Profiling](../tilelang/profiling-tilelang-programs/SKILL.md) |
| Which bounded schedule change tests that finding? | [Optimization](../tilelang/optimizing-tilelang-programs/SKILL.md) |
| How should the supported design be explained? | [Design documents](../tilelang/writing-tilelang-design-docs/SKILL.md) |

Mark a route unused when it does not answer the current question. Record an
unavailable specialist/tool and the resulting evidence gap; do not substitute
a prose guess for its check. Native CUDA compilation or sanitizer work can use
the [CUDA workflow](../cuda_skill/SKILL.md) when that boundary is relevant.

## 1. Establish the contract and baseline

Read the reference, current candidate, caller, and harness. Record the following
in the existing task record, or in `docs/draft.md`:

- Operation and input/output structure, shapes, dtypes, strides, accuracy,
  mutation/aliasing, initialization, stream, and gradient requirements.