tilelang-op-develop

name	tilelang-op-develop
description	根据冻结的 DESIGN.md 生成算子实现（{op}.py：kernel + golden），执行测试并返回三态判定。触发：实现算子、生成 kernel、算子开发、跑精度。

TileLang-NPUIR 算子开发与验证

1. 目标

根据 Stage 1 冻结的 DESIGN.md 与 Stage 2 通过的 REVIEW.md，生成算子实现文件 {op}.py（含 @tilelang.jit kernel + 内嵌 PyTorch golden + main 入口），执行测试，并返回三态判定供 conductor 路由。

环境前提：本 skill 运行在已具备 NPU 设备的环境中，tilelang 与 torch_npu 可正常导入。kernel 编译与执行在 NPU 上真实进行，精度校验为真实结果。

---

2. 输入

字段	说明
design_md_path	冻结的 DESIGN.md（含 L0 测试计划）
review_md_path	Stage 2 通过的 REVIEW.md（设计已检视通过）
mode	first_impl / retry_impl / precision_fix（由 conductor 传入）
attempt_index	当前 Stage 3 attempt 序号
last_failure_summary	重试时传入的失败信息（stderr 摘要 / 精度失败详情）
design_revision_count	设计修订次数（用于回退后清零判断）

---

3. 工作流程

Phase 1：读取设计

1. Read DESIGN.md 全文，提取：算子名、I/O 规格、编程模式、API 映射、Tiling、内存层级、同步策略、L0 测试计划、精度标准。
2. Read REVIEW.md，确认检视已通过（如有 warn 项记录但不阻塞）。

Phase 2：生成 kernel

1. 按 DESIGN.md §3 API 映射 + §6 循环结构生成 @tilelang.jit(target="npuir") kernel。
2. 优先 v-prefix API（vadd/vmul/vexp/vcast/vbrc），npuir_xxx 仅作兼容。
3. 遵循项目根 AGENTS.md："不要凭记忆猜 API"、"从示例入手"——先 Glob examples/ 同类实现参考。

Phase 3：生成 golden

1. 按 DESIGN.md §8.1 生成 PyTorch CPU 参考实现 golden_{op}(...)。
2. Golden 必须在 CPU 上可独立运行（不依赖 torch_npu）。

Phase 4：执行测试

1. 跑 L0：python {op}.py --level L0。
2. L0 通过后扩展 L1/L2/Boundary 并跑全量 --level all。
3. 收集结果：max_diff、失败用例 shape、层级。

Phase 5：三态判定与返回

条件	返回标记
输出与 golden 函数输出对比精度正常	[PRECISION_PASS]
输出与 golden 函数输出对比精度未过	[PRECISION_FAIL]
发现设计层错误（API 不可用、L0C 溢出、内存层级冲突等实现层无法修复）	[DESIGN_ERROR] + 原因
无标记且 exit code ≠ 0	运行失败（conductor 按 retry_impl 路由）

---

4. {op}.py 结构规范

生成的文件必须包含以下组成部分（顺序）： 注意：*.py 中的注释只能使用英文

# 1. Copyright (c) Huawei Technologies Co., Ltd. 2026.
# 2. imports (tilelang or torch_npu)
# 2. golden_{op}(...) function
# 3. @tilelang.jit kernel
# 4. precision comparing func：run_case()
# 5. main()

完整可运行模板见 templates/op_template.py.

main 块结构

def main():
    parser = argparse.ArgumentParser()
    parser.add_argument("--level", default="L0", choices=["L0", "all"])
    args, _ = parser.parse_known_args()

    if args.level == "L0":
        run_L0()
    else:
        run_L0(); run_L1(); run_L2(); run_boundary()

每个 run_LX() 内部在 NPU 上创建张量、调用 kernel 执行、与 golden 对比精度。详见模板。

---

5. 失败处理

失败类型	识别	处理
编译错误（实现层）	stderr 含 lowering/codegen 错误	返回运行失败 + stderr 摘要，conductor 走 retry_impl
API 不存在	AttributeError / 设计用 API 无导出	返回 [DESIGN_ERROR] + 原因
L0C/UB 溢出	编译期或运行期报容量超限	返回 [DESIGN_ERROR] + 原因
精度不达标	assert_close 失败	返回 [PRECISION_FAIL] + max_diff/失败 shape
内存层级越级	stderr 提示 GM/L1/UB/L0 访问违规	返回 [DESIGN_ERROR] + 原因
环境问题	ImportError 指向 tilelang/torch_npu 未安装或未 source set_env.sh	返回运行失败，提示检查环境

---

6. 备份规则

precision_fix 模式每次修改 {op}.py 前，必须先备份：

cp {op}.py history_version/{op}_impl_s3_attempt{N}.py

（{N} = 当前 attempt_index）

---

7. 完成报告

返回结构化摘要：

## Stage Result
- stage: 3
- mode: first_impl / retry_impl / precision_fix
- project: {project}
- operator: {op}
- output: examples/{project}/{op}/{op}.py
- verdict: [PRECISION_PASS] / [PRECISION_FAIL] / [DESIGN_ERROR] / RUNTIME_FAIL
- test_results:
  - L0: pass / fail (N cases)
  - L1: pass / fail (N cases)
  - L2: pass / warn (N cases, 不阻塞)
  - Boundary: pass / warn (N cases, 不阻塞)
- max_diff: <精度数值>
- design_error_summary: <仅 DESIGN_ERROR 时填>
- skills_consulted: <引用的 skill 路径>
- summary: <一句话>
- issues: <若无则 none>