tilelang-ascend: Skill for Claude Code

TileLang Ascend Tile API 开发流程

当任务是新增一个面向用户的 T.tile.xxx 小 API 时使用本 skill，尤其是从 tilelang/language/ascend_tile.py 暴露、并由 Ascend C 代码生成支撑的 API。

这里的目标不只是增加一个 Python 函数名，而是交付一个真正可用的 API：它能编译、能 lowering、能生成合法的 Ascend C、有清晰稳定的语义边界，并且有合适的 CI 测试覆盖。

第一轮调研

开始实现前，先读清楚当前仓库形态，再决定方案：

1. 阅读 AGENTS.md。
2. 阅读 .agents/skills/tilelang-custom-skill/tilelang-api-best-practices/SKILL.md。
3. 阅读 .agents/skills/tilelang-custom-skill/tilelang-api-best-practices/references/api-compute.md。
4. 如果涉及编程模式或 pass_configs，阅读 .agents/skills/tilelang-custom-skill/tilelang-programming-model-guide/SKILL.md。
5. 查看 tilelang/language/ascend_tile.py 中最相近的现有 API。
6. 查看 testing/python/language/ 中最相近的测试。
7. 查看 src/op/ascend.{h,cc}、src/target/codegen_ascend.cc、src/tl_templates/ascend/common.h 中相近的 lowering、codegen 和 helper 实现。

不要凭记忆推断 API 签名。已有本地模式优先于看起来更聪明的新抽象。

范围决策

编辑前，先明确并说明 API 边界：

• 用户 API 名称，例如 T.tile.foo(dst, src, ...)。
• 它属于纯 tile 计算、数据搬运、类 reduction 行为，还是带副作用的写回。
• 支持的 buffer scope，通常是 GM、UB（shared.ub）、L1（shared.l1）、L0（wmma.*），或其中子集。
• 支持的 dtype 和 rank。
• 是否接受 Buffer、BufferLoad、BufferRegion。
• 不支持的参数组合和语义。
• 应该支持 Developer pass_configs、Expert 模式，还是两者都支持。

当语义无法和主仓 / GPU 的全局 API 对齐时，优先新增 Ascend 专属的 T.tile.xxx API。除非用户明确要求、且语义确实一致，否则不要新增或修改全局 T.xxx API。

实现路径

1. Python 前端

在 tilelang/language/ascend_tile.py 中新增用户入口。

优先复用该文件和 tilelang/language/copy_op.py 里的本地 helper 模式：

• 如果现有 API 会解析 let-bound value，新 API 也应保持一致。
• 只在语义明确时接受 Buffer、BufferLoad 或 BufferRegion。
• 当 C++ lowering 需要 region、rank 或 extent 信息时，把前端输入转换为 tl.region。
• 对不支持的 scope 或参数组合尽早报错，并给出清晰错误信息。
• 当操作需要 C++ lowering 时，发射命名清晰的 op，通常是 tl.ascend_<api_name>。

除非操作本质上就是普通 copy，否则不要为了省事复用 tl.ascend_copy。有副作用、会改变硬件模式、或语义不同的操作，通常应该有显式 op 名称。

2. C++ Operator 和 Lowering

如果前端发射新的 tl.ascend_* op，需要补齐对应 C++ operator 路径：

• 在 src/op/ascend.h 中声明。
• 在 src/op/ascend.cc 中实现参数解析和 lowering。
• 当它需要 tile-op lowering 时，用 TIR_REGISTER_TL_OP 注册。
• 保留 region 信息，直到 lowering 可以计算访问指针、extent、stride、mask 或合法 shape。
• 即使 Python 已经做过校验，C++ 侧也要重新检查关键约束。