tile-ai TileLang-Ascend 算子设计文档生成
1. 目标
根据算子需求信息，生成一份完整的 TileLang-Ascend 算子设计文档（design.md），涵盖以下核心决策：
- 编程模式选型：Developer / Expert / 混合模式
- API 映射：将数学公式拆解为 TileLang DSL 原语组合
- 内存层级规划：GM → L1/UB → L0 的数据搬运路径
- Tiling 策略：Block 划分与 Tile Shape 设计
- 循环结构：T.Parallel / T.serial / T.Pipelined / T.Persistent 的选择
- 同步策略：自动同步 vs 手动同步标志
- 验证方案：Golden 函数与 L0 门槛测试计划（完整分层套件 L1/L2/Boundary 由 tilelang-op-test-design 生成）
2. 输入要求
必需信息
| 字段 | 说明 |
|---|---|
| 算子名称 | 如 softmax、layer_norm、flash_attention |
| 数学公式 | 算子的数学表达，如 $\text{softmax}(x_i) = e^{x_i} / \sum e^{x_j}$ |
| 输入张量规格 | shape、dtype |
| 输出张量规格 | shape、dtype |
| 编程模式偏好 | Developer / Expert / 混合 |
| 迁移算子路径 ⭐ | 原算子文件路径（迁移时必需），用于获取 golden 实现 |
| 输出形状 ⭐ | 原算子输出 shape（迁移时必需），如 (N, M) 或 (M, N) |
迁移算子时必须提供原算子路径和输出形状，否则无法证明迁移正确性。Golden 实现一致性要求详见 tilelang-op-develop checklist.md #9 Golden 实现一致 / #10 输出形状匹配。
提问规则（必须严格遵守）：
- 优先使用调用方传入的字段：若调用方（如
@tilelang-op-orchestrator通过 analyst 传入op_requirements结构）已经提供了字段值，全部跳过提问，直接进入技术约束检测和 design 生成 - 每次只询问一个字段：使用
question工具时，questions数组中只包含一个元素 - 按表格顺序依次询问：算子名称 → 数学公式 → 输入张量规格 → 输出张量规格 → 编程模式偏好
- 已提供的字段跳过：如果用户在初始请求中已提供某个字段的值，跳过该字段继续下一个
- 示例：
- 第 1 次询问：只问"数学公式"
- 用户回答后，第 2 次询问：只问"输入张量规格"
- 以此类推
⚠️ 当被 orchestrator → analyst Subagent 链路调度时：
- analyst 会把 orchestrator 在 Primary 上下文预检收集到的
op_requirements完整传入 - 此时 5 个必需字段应当全部已 provided，跳过整个提问环节
- 若 skill 仍发现字段歧义或缺漏，不要在当前 Subagent 上下文调用
AskUserQuestion（透传不到真实用户），而是让 analyst 返回partial_input+ 缺失字段名给 orchestrator，由 orchestrator 在 Primary 上下文追问
推荐信息
| 字段 | 说明 |
|---|---|
| 典型配置 | 常用的 shape 组合与优先级 |
| 参考实现 | PyTorch / NumPy 参考代码 |
| 性能目标 | 目标吞吐量或延迟 |
| 动态轴说明 | 哪些维度在运行时变化 |
若用户未提供必需信息中的任一项，通过提问补全后再继续。
3. 技术约束（必须遵守）
本项目为 TileLang-Ascend（华为昇腾 NPU），与 GPU 版 TileLang 有显著差异。外部参考实现不可直接使用，必须转换为 Ascend 兼容方案。
生成 design.md 前必须执行强制检测：三维 Kernel、threads 参数、动态循环边界、GPU 专用 API、GEMM 非整除、L0C 溢出等。
详细已知限制清单、强制检测规则、警告输出模板见 references/ascend-constraints.md。
4. 工作流程
Phase 1：输入解析与算子特征分析
- 解析算子名称与数学公式
- 验证必需字段是否完整
- 分析算子特征：
- 计算类型判定：
- 纯 Vector（element-wise / reduction）→ 仅需 UB
- 纯 Cube（仅 matmul）→ 需要 L1 + L0A/L0B/L0C
- 混合（matmul + element-wise 后处理）→ 核间流水线，需要 CV 融合
- Host 预处理：如 im2col 等 Python 侧预处理步骤，标明在 design 的 §1 和 §4 中
- 复杂度级别：
- 单步（如 element-wise add）→ 无循环、单次搬运
- 多步（如 softmax = max + sub + exp + sum + div）→ 多次计算、可能需要中间缓冲
- 融合（如 flash attention = GEMM + softmax + GEMM）→ 核间协作、流水线
- 动态 shape 判定：是否存在运行时才确定的维度
- 计算类型判定：
- 非整除场景预判：检查输入 shape 是否可能不被 block size 整除。
T.ceildiv(M, block_M)对非整除或M < block_M返回 ≥1（非零），T.copy已支持动态 shape 切片自动处理尾块，不需要 host padding。用T.ceildiv+ 动态切片T.copy(A[m:m+valid, ...])，参考examples/chunk_gated_delta_rule/expert_chunk_gated_delta_rule.py:107-108。仅当多个 group 共享同一输出 buffer 时需注意尾块写入竞态——用 metadata 的 valid_m 字段限制写入范围T.copy(C_L0, Y[m:m+valid_m, ...]) - 多 group 输出竞态约束（grouped 类算子）：当多个 group 共享同一输出 buffer（紧凑排列，不 padding）时，尾块按 block_M 整块写会溢出到隔壁 group 的区域，导致竞态条件（执行顺序不确定→结果不确定）。解法：metadata 记录 valid_m，kernel 用
T.copy(C_L0, Y[m_start : m_start + valid_m, ...])只写有效行。参考examples/grouped_gemm/example_grouped_gemm_fwd.py的 block_metadata[2]（valid_m 字段，当前未使用，应启用）。 - 数据重排成本建模：按 references/ascend-constraints.md §4.2 为每条路径和最大/关键 case 计算 DMA/GM 标量访问/地址解码/并行度；不能只给 tile shape，不计算 transaction 数量。
Phase 2：信息收集
必须执行强制步骤 0：搜索本项目同类实现。详细工具调用、信息收集步骤、禁止行为见 references/info-sources.md。
Phase 3：生成 design.md
⚠️ 生成 design.md 时必须遵守 references/ascend-constraints.md §5「Host 侧 Buffer 操作约束」。下游
tilelang-op-develop会再次校验；违规设计在 Stage 2 返回[DESIGN_ERROR]。
基于 examples/design-template.md 模板，填充所有章节：
- 概述
- 编程模式选型
- API 映射设计
- 数据规格与内存规划
- Tiling 策略（非整除时用输入、输出两侧显式 valid extent/BufferRegion；前端负责 按这些动态切片裁剪搬运，不代表尾块无需设计；host 侧不允许 padding + crop）
- 循环与调度结构
- 同步策略
- CV 融合设计（按模式分支：Developer 默认消除 workspace/vid——
threads=2+ 片上直连，不产出 workspace 规格；仅 Expert/混合或复杂场景回退才设计 workspace +workspace_idx。详见 design-template.md §8.2） - 验证方案（Golden + L0 门槛测试计划 + 性能可行性哨兵；除规则
shape 外，至少包含每条路径最坏 dtype/最大任务数的用户关键 case，并给出单 case 超时预算；
完整分层套件 L1/L2/Boundary 交由
tilelang-op-test-design）。若算子支持 NaN/Inf，精度方案必须声明位置敏感比较：特殊值用“有限值 + 稀疏特殊值”的 混合输入，先严格比较 NaN/正 Inf/负 Inf mask，再只对有限值应用数值容差； 禁止使用全 NaN/全 Inf 输入作为唯一特殊值用例，因为数据重排或索引错误可能被掩盖 - 风险点与注意事项
- 交付清单
Phase 4：质量自检
⚠️ 首要检查：host 侧 Buffer 操作合规性（违反则立即修订，不得继续）
核对 design.md 的完整 host 路径是否只含经证明不物理化的 metadata view、kernel 调用与验证；命中真实拷贝或 aclnn 调用必须修订（详见 references/ascend-constraints.md §5）。
按照 references/quality-checklist.md 中的自检清单逐项检查，确保文档质量。
Phase 5：针对性修订
仅修正未通过自检的项目。信息确实不足的标注为「待确认」并说明原因。
Phase 6：输出
- 将
design.md输出到当前目录或用户指定路径。若文件已存在，询问是否覆盖。 - 同时产出
proto.yaml（算子接口规格，模板见 examples/design-template.md §11.5）：dtype 全集取自 §9.3 精度表（每个支持的 dtype 一行；§4.1 只给代表性 dtype，不作 dtype 全集来源）、attr 取自 §1/§4，机械派生写到同目录（examples/{op}/proto.yaml）。这是覆盖门禁coverage_check.py --proto的权威 dtype/attr 来源，每个算子都必须产出；inputs[].dtype须与 §9.3 精度表的 dtype 行一致。
5. 算子特征分析决策树
详细决策树（Ascend 版）、平台识别、API 映射规则、NPU 硬件约束（分形限制 / 对齐要求 / 存储大小上限）见 references/decision-tree.md。
6. 信息源优先级
信息源优先级表与冲突处理原则见 references/info-sources.md。
7. 错误处理
| 场景 | 处理方式 |
|---|---|
| 用户未提供数学公式 | 提问补全，给出常见算子公式作为参考 |
| 必需字段缺失 | 列出缺失项，逐一提问 |
| API 查询无结果 | 标注为「需扩展」，在风险点中说明 |
| 目标文件已存在 | 询问用户是否覆盖或另存 |
| 算子过于复杂 | 建议拆分为多个子算子分别设计 |
8. 完成报告
文档生成完成后，按 examples/completion-report-template.md 输出报告。
9. 生成算子
完成报告后，询问用户是否根据此报告生成对应算子代码。
子目录索引
- references/ascend-constraints.md — 技术约束清单、强制检测规则、警告输出格式
- references/decision-tree.md — 算子特征分析决策树、平台识别、NPU 硬件约束、API 映射规则
- references/quality-checklist.md — 质量自检清单
- references/info-sources.md — 信息收集步骤、信息源优先级、冲突处理原则
- examples/design-template.md — design.md 完整模板
- examples/completion-report-template.md — 完成报告输出模板