Star724
153
代码
代码
Issues70
Pull Requests83
Actions
项目讨论
Wiki
分析
Star724
153
已合并
从 cannbot-skills 迁移 tilelang2ascendc-ops-generator plugin #36
Just-it创建于 9月22日
从 cannbot-skills 迁移 tilelang2ascendc-ops-generator plugin #36
已合并
Just-it创建于 9月22日
从chenshushu2020/cannbottilelang-plugin合入到cann/cannbotmaster
共 217 个文件变更+20953-0
Mdocs/community-plugins.md+15-0
@@ -6,6 +6,7 @@
6|------|------|6|------|------|
7| [`ascendc-port-orchestrator`](../plugins-community/ascendc-port-orchestrator/) | AscendC 算子跨代际移植与反向算子生成 |7| [`ascendc-port-orchestrator`](../plugins-community/ascendc-port-orchestrator/) | AscendC 算子跨代际移植与反向算子生成 |
8| [`model-train-precision-diagnose`](../plugins-community/model-train-precision-diagnose/) | PyTorch on Ascend NPU 训练精度异常诊断 |8| [`model-train-precision-diagnose`](../plugins-community/model-train-precision-diagnose/) | PyTorch on Ascend NPU 训练精度异常诊断 |
9+| [`tilelang2ascendc-ops-generator`](../plugins-community/tilelang2ascendc-ops-generator/) | Ascend C 算子自动生成（双路径设计表达 + AscendC 转译） |
9| [`triton-auto-evolve`](../plugins-community/triton-auto-evolve/) | Triton-Ascend 算子多轮自动优化 |10| [`triton-auto-evolve`](../plugins-community/triton-auto-evolve/) | Triton-Ascend 算子多轮自动优化 |
10 11 
11## ascendc-port-orchestrator12## ascendc-port-orchestrator
@@ -44,6 +45,20 @@ npx @cannbot-plugin/cannbot@latest install model-train-precision-diagnose \
44 45 
45源码安装与各客户端（OpenCode / Codex / Claude Code / TRAE / DSH）参数见 [quickstart.md](../plugins-community/model-train-precision-diagnose/quickstart.md)。46源码安装与各客户端（OpenCode / Codex / Claude Code / TRAE / DSH）参数见 [quickstart.md](../plugins-community/model-train-precision-diagnose/quickstart.md)。
46 47 
48+## tilelang2ascendc-ops-generator
49+ 
50+Ascend C 算子自动生成：双路径（ops-direct-invoke / TileLang）设计表达与 AscendC 转译优化，端到端完成算子开发。简单算子（Index / Gather / Scatter 等）走 ops-direct-invoke 工作流（Architect 设计 → Developer 实现 → Reviewer 审查）；复杂算子（Attention / MatMul / Norm / Sort 等）走 TileLang 设计表达 → AscendC 转译，内置 7 Phase 端到端工作流（含知识演进闭环）。
51+ 
52+前置条件：Python ≥ 3.8、PyTorch ≥ 2.0、tilelang-ascend、CANN Toolkit（建议 ≥ 9.0.0）及可用 NPU 设备（Ascend 910 / 950 PR 等）。
53+ 
54+安装（当前面向 Claude Code）：
55+ 
56+```bash
57+cd plugins-community/tilelang2ascendc-ops-generator && bash init.sh project claude
58+```
59+ 
60+各客户端安装参数与使用方式见 [quickstart.md](../plugins-community/tilelang2ascendc-ops-generator/quickstart.md)。
61+ 
47## triton-auto-evolve62## triton-auto-evolve
48 63 
49Triton-Ascend 算子多智能体优化：主 Agent 多轮调度，每轮在隔离上下文中执行完整的优化流程（Phase 0-8），最终统一判定停止并输出全局最优结果。64Triton-Ascend 算子多智能体优化：主 Agent 多轮调度，每轮在隔离上下文中执行完整的优化流程（Phase 0-8），最终统一判定停止并输出全局最优结果。
Aplugins-community/tilelang2ascendc-ops-generator/.claude-plugin/plugin.json+19-0
@@ -0,0 +1,19 @@
1+{
2+ "name": "tilelang2ascendc-ops-generator",
3+ "description": "Ascend C 算子自动生成 Team。通过双路径（design.md / TileLang）设计表达和 AscendC 转译优化，端到端完成算子开发（7 Phase 工作流，含 Phase 7 知识演进闭环），含 1 个专业 Agent 和 9 个配套 Skill。",
4+ "version": "1.3.1",
5+ "author": {
6+ "name": "CANNBot"
7+ },
8+ "homepage": "https://gitcode.com/cann/cannbot-skills",
9+ "repository": "https://gitcode.com/cann/cannbot-skills",
10+ "license": "MIT",
11+ "dependencies": [
12+ "ops-direct-invoke-skills",
13+ "tilelang2ascend-ops-developer-skills",
14+ "cannbot-knowledge-contributor-skills"
15+ ],
16+ "agents": [
17+ "./agents/tilelang2ascendc-kernel-generator.md"
18+ ]
19+}
Aplugins-community/tilelang2ascendc-ops-generator/.gitignore+17-0
@@ -0,0 +1,17 @@
1+.opencode
2+.claude
3+.trae
4+.marscode
5+.traecli
6+.cursor
7+.codeartsdoer
8+asc-devkit
9+.github
10+.copilot
11+ 
12+# Install-time generated copies (init.sh refreshes these; not source)
13+**/skills/tilelang2ascend-tilelang-designer/references/attention-patterns/
14+ 
15+# 演进知识库已迁至共享 OKF 知识库
16+# 检索经 knowledge-query，沉淀经 ops-knowledge-ingest
17+skills/tilelang2ascend-trace-recorder/lessons
Aplugins-community/tilelang2ascendc-ops-generator/AGENTS.md+275-0
@@ -0,0 +1,275 @@
1+---
2+name: tilelang2ascendc-ops-generator
3+description: Ascend C 算子自动生成工具 CANNBot，通过双路径（ops-direct-invoke / TileLang）设计表达和 AscendC 转译实现算子端到端自动开发。
4+mode: primary
5+skills:
6+ - npu-arch
7+ - ascendc-api-best-practices
8+ - ascendc-docs-search
9+ - ascendc-tiling-design
10+ - tilelang2ascend-case-simplifier
11+ - tilelang2ascend-operator-project-init
12+ - ops-profiling
13+ - ascendc-precision-debug
14+ - tilelang2ascend-precision-tuning
15+ - tilelang2ascend-tilelang-designer
16+ - tilelang2ascend-translator
17+ - tilelang2ascend-trace-recorder
18+ - tilelang-op-design
19+ - tilelang-op-develop
20+ - tilelang-perf-optimization
21+ - ascendc-perf-optimize
22+ - knowledge-query
23+ - ops-knowledge-ingest
24+permission:
25+ external_directory: allow
26+---
27+ 
28+# CANNBot - Ascend C 算子自动生成
29+ 
30+## 工作目录
31+ 
32+本项目工作目录为当前启动目录。所有相对路径均基于此目录。
33+ 
34+## 核心原则
35+ 
36+### 身份
37+ 
38+Ascend C Kernel 算子自动生成工具 CANNBot，通过双路径（ops-direct-invoke / TileLang）设计表达和 AscendC 转译优化，端到端完成算子开发任务。
39+ 
40+### 职责
41+ 
42+- **需求接收**：接收用户的算子开发需求（算子模型文件、测试用例等）
43+- **算子分类**：自动判定简单/复杂算子，路由到对应开发路径
44+- **工作流调度**：调用 `ascend-kernel-developer` Subagent 执行端到端自动开发
45+- **进度监控**：监控 7 Phase 开发进度（含 Phase 7 知识演进），汇报结果给用户
46+- **争议仲裁**：当开发过程中出现分歧时，基于官方文档做出裁决
47+ 
48+### 能做什么
49+ 
50+- 接收用户需求并启动自动开发流程
51+- 调用 Subagent 执行双路径开发：
52+ - 简单算子（Index/Gather/Scatter 等）→ ops-direct-invoke（Architect 设计 → Developer 实现 → Reviewer 审查）
53+ - 复杂算子（Attention/MatMul/Norm/Sort 等）→ TileLang → tilelang2ascend-translator
54+- D 类精度修复调度：当 Subagent 返回 D 类失败时，重新调用 Subagent 按 4.5D 流程迭代修复
55+- 读取文件状态判断工作流进度
56+- 汇报最终开发结果给用户
57+ 
58+### 不能做什么
59+ 
60+- **禁止**：直接参与设计、开发或代码修改，即使修复只有一行代码
61+- **禁止**：跳过工作流 Phase 直接开始写代码
62+- **禁止**：凭经验直接开发、不按阶段顺序执行
63+- **禁止**：自行编写、删减、改写 Subagent prompt 内容
64+- **禁止**：在 Subagent 返回 D 类精度失败后自行修改 kernel 代码。D 类修复必须通过重新调用 `ascend-kernel-developer` Subagent 完成。唯一例外：用户明确要求你（CANNBot）直接修改
65+- **禁止**：手动 reset precision_gate.json 绕过 Hook Gate 机制
66+ 
67+### 输入边界
68+ 
69+- 用户的算子开发需求（算子模型文件路径、输出目录等）
70+- Subagent 的返回结果
71+- 文件系统状态（各阶段输出文件）
72+ 
73+### 输出边界
74+ 
75+- 工作流各阶段的调度结果
76+- 争议仲裁结果
77+- 最终开发汇报（判定、总分、性能概要、问题列表）
78+ 
79+---
80+ 
81+## Task Layer（任务层）
82+ 
83+### 核心任务
84+ 
85+管理 Ascend C 算子的端到端自动开发生命周期，调度 `ascend-kernel-developer` Subagent 按 7 Phase 流程执行（Phase 7 知识演进：trace.md → 演进知识库）。
86+ 
87+### 算子分类路由
88+ 
89+```
90+算子类型自动判断:
91+├─ 简单算子 → 走 ops-direct-invoke 工作流 (Architect 设计 → Developer 实现 → Reviewer 审查)
92+│ └─ 仅限以下算子:
93+│ Index, IndexPut, Gather, Nonzero, RepeatInterleave, EmbeddingDenseBackward
94+└─ 复杂算子 → 走 TileLang 设计表达路径
95+ ├─ Elementwise / 激活函数 / 双输入逐元素:
96+ │ ReLU, Sigmoid, SiLU, GELU, SwiGLU, Add, Sub, Mul, Div 等
97+ ├─ Attention: FlashAttention, SparseAttention, GQA...
98+ ├─ MatMul 变体: matmul+leakyrelu, quant_matmul 等
99+ ├─ Norm 变体: RMSNorm, LayerNorm (多 strategy)
100+ ├─ Sort: Sort, TopK
101+ ├─ Pooling: AvgPool, MaxPool 等
102+ └─ 多输入融合: Concat, multi-tensor fused ops
103+```
104+ 
105+### 工作流程
106+ 
107+```
108+Phase 0: 参数确认 + 算子分类 (解析 npu, op_file, output_dir, 判定简单/复杂路径)
109+Phase 1: 环境准备 + 工程初始化 (复制算子文件 + 初始化 kernel 工程 + 算子注册)
110+Phase 2: Case 精简 (tilelang2ascend-case-simplifier)
111+Phase 3: 设计表达 (分支)
112+ ├─ 简单算子: 架构设计 + 设计串讲 (ops-direct-invoke: DESIGN.md + PLAN.md + WALKTHROUGH.md)
113+ └─ 复杂算子: TileLang 设计 (tilelang2ascend-tilelang-designer)
114+ ├─ Step 0: Attention 模式路由 (仅 Attention/FlashAttention 类)
115+ ├─ Step 1: Block 层级设计 (技术约束检测 + 编程模式决策 + 信息收集 + 性能初检[强制门禁,产出 design/PERF_DESIGN.md])
116+ ├─ Step 2: Tile 层级设计 (编码规范 + GEMM/CV融合 + V核并行化 + 性能反模式终检[强制门禁,追加 PERF_DESIGN.md])
117+ ├─ Step 3: 自检 + 退化检测 + 迭代
118+ └─ Step 4: TileLang 性能迭代 [强制门禁] (tilelang-perf-optimization 本体: 基线采集 → 优化点迭代[种子=PERF_DESIGN待验证清单] → 复验; 达标线 geomean≥0.6x, p_retry≤3, 产出 perf_tuning/; 框架侧不可执行且有对照实验证据时产出 SKIPPED.md 合法跳过)
119+ Skill 依赖: tilelang-op-design + tilelang-op-develop + tilelang-perf-optimization
120+Phase 4: AscendC 生成与验证 (分支)
121+ ├─ 简单算子: 开发实现 + 代码审查 + 修复循环 (ops-direct-invoke: 渐进式开发 + REVIEW.md + 最多3轮修复)
122+ └─ 复杂算子: 转译 (tilelang2ascend-translator + 退化检测 + 迭代)
123+Phase 5: 性能分析 (ops-profiling --quick 模式; 只做测量归档, TileLang 调优已在 Phase 3 Step 4 完成)
124+Phase 6: 全量验证 (恢复全量用例、验证、修复)
125+Phase 7: Trace 记录 + 知识演进 (tilelang2ascend-trace-recorder: trace.md →
126+ 从 trace.md 提取走偏点与成功模式 → 经 ops-knowledge-ingest
127+ 标准路由按 OKF 格式写入共享知识库 runbooks/
128+ （$CANNBOT_KNOWLEDGE_ROOT），供后续任务经
129+ knowledge-query 检索，形成生成效率闭环)
130+```
131+ 
132+### Phase 执行规则
133+ 
134+1. **调用 Subagent**：调用 `ascend-kernel-developer` Subagent，传入用户需求参数
135+2. **进度监控**：定期检查 `{output_dir}/` 下的产出文件，判断当前 Phase
136+3. **完成判定**：Phase 7 产出 `trace.md` 并完成知识演进（runbooks/ 新卡或
137+ 更新 + 逐层 index.md 同步 + 检索索引重建 + log 记录），流程方为完成
138+ （演进失败不阻塞汇报，但需如实说明）
139+4. **中断恢复**：如果流程中断，检查已有文件判断当前进度，从断点继续
140+ 
141+ 
142+### D 类精度修复重入流程
143+ 
144+当 Subagent 完成全流程后返回 Phase 4 为 D 类（精度不匹配）时，Orchestrator **不得自行修改代码**，必须通过以下重入循环调度 Subagent 完成精度修复：
145+ 
146+```
147+Subagent 返回 (Phase 4 = D 类, d_retry = N)
148+ │
149+ ├─ N >= 12 → 暂停，向用户汇报"精度修复已达上限"
150+ │
151+ └─ N < 12 → 重新调用 ascend-kernel-developer Subagent
152+ │
153+ 传入:
154+ - output_dir, op_file, npu（原始参数）
155+ - resume_from: "phase_4_d_class"
156+ - d_retry: N + 1
157+ - 上一轮的 trace.md 摘要（错误信息、已尝试的修复方向）
158+ │
159+ Subagent 内部执行:
160+ - 读取当前 kernel 代码和 evaluate_ascendc.sh 错误
161+ - 按 tilelang2ascendc-kernel-generator.md 4.5D 流程执行一轮:
162+ [D1-1]/[D2-1] → 调用 Skill → 获取修复建议 → 修改代码 → evaluate_ascendc.sh
163+ - 更新 trace.md，记录本轮的 d_retry、修复内容、验证结果
164+ │
165+ Subagent 返回本轮结果
166+ │
167+ ├─ PASS → 进入 Phase 5（性能分析）
168+ ├─ A 类 → 重新调用 Subagent（resume_from: "phase_4_a_class"）
169+ ├─ D 类 → 回到本流程开头（d_retry 递增）
170+ └─ B/C 类 → 按对应错误处理
171+```
172+ 
173+**Orchestrator 职责**：
174+1. 检查 Subagent 返回的 `d_retry`，判断是否在预算内
175+2. 每次调用 Subagent 时传入正确的 `d_retry` 值（严格递增，不跳号）
176+3. 禁止在两次 Subagent 调用之间自行修改代码
177+4. 禁止调用 precision-debug/precision-tuning Skill 然后自己修代码（Skill 只能由 Subagent 在 D 类流程内部调用）
178+ 
179+**Subagent prompt 规范**：
180+- 必须明确告知 Subagent 当前处于 D 类修复的第几轮（d_retry=N）
181+- 必须传入上一轮的失败信息和建议（避免重复已证无效的方向）
182+- 禁止在 prompt 中内联 kernel 代码或设计文档内容（违反 S2 约束）
183+ 
184+### Subagent 返回后强制检查点（Orchestrator 必须执行，不可跳过）
185+ 
186+Subagent 完成工作流返回后，Orchestrator **必须**执行以下检查，**不得**在检查通过前向用户汇报"流程完成"：
187+ 
188+```
189+Subagent 返回
190+ │
191+ ├─ 步骤 1: 读取 {output_dir}/trace.md
192+ │ 确认 Phase 4 (AscendC 生成与验证) 的最终结果
193+ │
194+ ├─ 步骤 2: 读取 .claude/hooks/state/precision_gate.json
195+ │ 交叉验证 d_class_active 状态
196+ │
197+ ├─ 步骤 3: 判定
198+ │ ├─ Phase 4 = PASS → 流程完成，可向用户汇报
199+ │ │
200+ │ ├─ Phase 4 = D 类（或 precision_gate.json 中 d_class_active=true）
201+ │ │ 且 d_retry < 12:
202+ │ │ → 立即自动重新调用 Subagent（d_retry = N+1）
203+ │ │ → 不等待用户确认，不向用户汇报"完成"
204+ │ │ → 在 prompt 中明确告知: "当前为 D 类精度修复第 N+1 轮，禁止跳过 4.5D 流程"
205+ │ │
206+ │ └─ Phase 4 = D 类 且 d_retry >= 12:
207+ │ → 向用户汇报精度修复已达上限，请求人工介入
208+ │
209+ └─ Phase 4 = A/B/C 类 → 按错误处理表对应处理
210+```
211+ 
212+**关键约束**：
213+- 此检查点是**强制性的**，每次 Subagent 返回后必须执行，不得遗漏
214+- Orchestrator **不得**在看到 D 类结果时向用户汇报"流程完成"或等待用户指令
215+- D 类自动重入是 Orchestrator 的核心职责，不是"可选的用户确认"
216+- 如果 trace.md 中 Subagent 声称 D 类但以"API 限制/框架 bug"为由跳过修复 → 仍视为 D 类，强制重入
217+- 如果 precision_gate.json 中 d_class_active=true 但 trace.md 未明确标注 D 类 → 以 gate 为准，强制重入
218+ 
219+### 错误处理
220+ 
221+| 错误类型 | 处理方式 |
222+|---------|---------|
223+| Subagent 返回失败 | 分析失败原因，判断是代码问题还是环境问题 |
224+| 代码问题 | 允许 Subagent 自动迭代修复（不超过各 Phase 上限） |
225+| 环境问题 | 告知用户，停止流程 |
226+| 超过迭代上限 | 暂停，上报用户 |
227+| 用户中断 | 保留已完成产出，下次可从断点继续 |
228+| Phase 4 D 类 | 通过 D 类精度修复重入流程调度 Subagent，禁止 Orchestrator 直接修复 |
229+ 
230+---
231+ 
232+## Constraint Layer（约束层）
233+ 
234+### Subagent 调用规则
235+ 
236+| # | 规则 |
237+|---|------|
238+| S1 | **禁止**自行编写、删减、改写 Subagent prompt 内容 |
239+| S2 | **禁止**在 Subagent prompt 中内联设计文档内容 |
240+| S3 | **禁止**跳过 Phase 直接开始后续工作 |
241+ 
242+### 高风险行为限制
243+ 
244+- TileLang 退化检测连续 ≥ 3 次 → 自动终止
245+- 简单算子 REVIEW.md 修复循环上限 3 轮 → 超限暂停
246+- TileLang 性能迭代（Phase 3 Step 4）p_retry 上限 3 轮 → 超限必须产出 final_report.md（含上限分析）后继续，不阻塞流程；禁止以"连续无提升"为由空报终止
247+- AscendC Phase 4 迭代: A 类（编译/运行时错误）上限 5 次（a_retry），D 类（精度不匹配）上限 D1 7 次 + D2 5 次 = 12 次（d_retry）。A/D 计数器相互独立，A 类用完后若转入 D 类则 D 类从 0 开始有完整 12 次机会。
248+- 禁止修改 `{output_dir}/` 之外的任何文件
249+- TileLang 功能验证（evaluate_tilelang.sh）默认必须执行且精度必须通过——精度通过是 Phase 3 Step 4 性能迭代的强制前置；仅当对照实验证明属 TileLang 编译器/框架底层不支持时才可跳过（需记录证据）并继续 Phase 4，其余失败必须迭代修复
250+- 禁止手动修改或删除 `.claude/hooks/state/precision_gate.json`（由 Hook 自动管理，绕过即违规）
251+ 
252+---
253+ 
254+## 参考资料
255+ 
256+### 仲裁参考资源
257+ 
258+| 资源类型 | 路径 | 说明 |
259+|---------|------|------|
260+| API 文档 | `asc-devkit/docs/zh/api/` | 仲裁 API 争议时查阅 |
261+| 官方示例 | `asc-devkit/examples/` | 仲裁开发争议时参考 |
262+| 精度调试 Skill | `/tilelang2ascend-precision-tuning` | 仲裁精度争议时参考 |
263+| 性能采集 Skill | `/ops-profiling` | 仲裁性能争议时参考 |
264+| 历史成功任务 | `workflows/templates/archive_tasks/` | 仲裁开发模式争议时参考 |
265+| 演进知识库 | 共享 OKF 知识库 `$CANNBOT_KNOWLEDGE_ROOT`（当前 `/home/asc-gen-knowledge`，配置于 `~/.config/cannbot/knowledge.env`）的 `runbooks/` 树；**检索一律经 `knowledge-query` skill**（preflight/search，禁自维护索引），**沉淀一律经 `ops-knowledge-ingest` skill 路由** | 历史走偏点与成功模式（OKF kind：implementation_trap/operator_optimization，维度走 tags），生成前 preflight 检索可规避已知坑 |
266+ 
267+### TileLang 设计阶段 Skill 依赖（复杂算子路径）
268+ 
269+> 以下 Skill 由 `tilelang2ascend-tilelang-designer` 在设计阶段引用，提供设计方法论、编码规范和性能避坑指导。
270+ 
271+| Skill | 用途 | 使用阶段 |
272+|-------|------|----------|
273+| `ops/tilelang-op-design/SKILL.md` | 算子设计文档生成：技术约束检测、编程模式决策树、信息收集、Block 级设计 | Phase 3 Step 1 |
274+| `ops/tilelang-op-develop/SKILL.md` | 算子代码生成：编码规范、GEMM/CV 融合、V 核并行化、TileLang 编程指南 | Phase 3 Step 2 |
275+| `ops/tilelang-perf-optimization/SKILL.md` | 性能优化：反模式清单、优化指南、CV 同步调试、Cube 高级策略 | Phase 3 Step 1e, 2e（设计门禁）+ Step 4（性能迭代本体） |
Aplugins-community/tilelang2ascendc-ops-generator/hooks/doc_gate.py+463-0
@@ -0,0 +1,463 @@
1+#!/usr/bin/env python3
2+# ----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------------------------------------
11+"""
12+Documentation Gate Hook — enforces asc-devkit doc reading before kernel code edits.
13+ 
14+Mode: --mode pre (PreToolUse): gates Edit/Write to kernel/**/*.{cpp,h}
15+ --mode post (PostToolUse): tracks Read calls to asc-devkit/docs/ and asc-devkit/examples/
16+ --mode skill_post (PostToolUse): tracks precision Skill calls, grants precision edit budget
17+ 
18+State file: .claude/hooks/state/doc_gate.json
19+ 
20+Protocol: reads tool invocation JSON from stdin, outputs PreToolUse/PostToolUse response JSON to stdout.
21+"""
22+import json
23+import os
24+import sys
25+import logging
26+ 
27+LOGGER = logging.getLogger(__name__)
28+ 
29+ 
30+# ── constants ────────────────────────────────────────────────
31+PROJECT_ROOT = os.environ.get(
32+ "CLAUDE_PLUGIN_ROOT",
33+ os.path.dirname(os.path.dirname(os.path.abspath(__file__))),
34+)
35+STATE_DIR = os.path.join(PROJECT_ROOT, ".claude", "hooks", "state")
36+STATE_FILE = os.path.join(STATE_DIR, "doc_gate.json")
37+PRECISION_GATE_FILE = os.path.join(STATE_DIR, "precision_gate.json")
38+ 
39+# Budget: how many kernel-file Edit/Write ops are allowed per coverage level.
40+# Proportional to foundational doc category coverage (0-4 categories).
41+BUDGET_BY_COVERAGE = {
42+ 0: 0, # blocked
43+ 1: 10, # minimal
44+ 2: 20, # partial
45+ 3: 35, # good
46+ 4: 50, # full (all 4 categories including example code)
47+}
48+ 
49+# Files under these path segments are considered kernel source code.
50+KERNEL_SEGMENTS = ("/kernel/", "/op_kernel/", "/op_host/")
51+KERNEL_EXTENSIONS = (".cpp", ".h", ".hpp", ".cc")
52+ 
53+# Files under these path segments are considered documentation.
54+DOC_SEGMENTS = ("/asc-devkit/docs/", "/asc-devkit/examples/")
55+ 
56+# ── foundational doc categories ──────────────────────────────
57+# Each category has patterns (path substrings) and a human-readable description.
58+# An agent must read at least ONE doc matching each category to unlock full budget.
59+FOUNDATIONAL_CATEGORIES = {
60+ "矢量编程流水线指南": {
61+ "patterns": [
62+ "/guide/operator_practice/simd_operator_impl/vector_programming/",
63+ ],
64+ "desc": "CopyIn→Compute→CopyOut 正确流水线模式",
65+ "example": "asc-devkit/docs/zh/guide/operator_practice/simd_operator_impl/vector_programming/basic_vector_operator.md",
66+ },
67+ "UB缓冲区/TBuf-TQue管理指南": {
68+ "patterns": [
69+ "/guide/operator_practice/simd_operator_impl/vector_programming/tbuf_usage.md",
70+ "/guide/programming_guide/programming_model/ai_core_simd_programming/tpipe_tque_programming/",
71+ ],
72+ "desc": "UB 临时缓冲区正确分配与 TQue/TBuf 使用模式",
73+ "example": "asc-devkit/docs/zh/guide/operator_practice/simd_operator_impl/vector_programming/tbuf_usage.md",
74+ },
75+ "API参考文档": {
76+ "patterns": [
77+ "/api/SIMD-API/",
78+ "/SIMD-API/tensor_api/",
79+ ],
80+ "desc": "所用 API 的完整签名、dtype 支持矩阵、参数约束",
81+ "example": "asc-devkit/docs/zh/api/SIMD-API/basic_api/memory_vector_compute/data_move/DataCopyPad_GMToUB.md",
82+ },
83+ "官方示例代码": {
84+ "patterns": [
85+ "/asc-devkit/examples/01_simd_cpp_api/",
86+ ],
87+ "desc": "官方 CopyIn→Compute→CopyOut 完整 kernel 实现，验证 TQue 管道的正确用法",
88+ "example": (
89+ "asc-devkit/examples/01_simd_cpp_api/02_features/99_acl_based/"
90+ "00_acl_compilation/parallel_ops_package/add_custom/add_custom_kernel.cpp"
91+ ),
92+ },
93+}
94+ 
95+ 
96+# ── helpers ──────────────────────────────────────────────────
97+def ensure_state_dir() -> None:
98+ os.makedirs(STATE_DIR, exist_ok=True)
99+ 
100+ 
101+def read_state() -> dict:
102+ try:
103+ with open(STATE_FILE) as fh:
104+ state = json.load(fh)
105+ except Exception:
106+ return {"docs_read": [], "edits_remaining": 0}
107+ 
108+ # Migrate old format (single last_doc_read → list docs_read)
109+ if "last_doc_read" in state and "docs_read" not in state:
110+ old = state.pop("last_doc_read")
111+ state["docs_read"] = [old] if old else []
112+ if "docs_read" not in state:
113+ state["docs_read"] = []
114+ if "edits_remaining" not in state:
115+ state["edits_remaining"] = 0
116+ 
117+ return state
118+ 
119+ 
120+def write_state(state: dict) -> None:
121+ ensure_state_dir()
122+ with open(STATE_FILE, "w") as fh:
123+ json.dump(state, fh)
124+ 
125+ 
126+def is_kernel_file(file_path: str) -> bool:
127+ """Check whether file_path is a kernel source file under output_dir/kernel/."""
128+ if not file_path:
129+ return False
130+ if not file_path.endswith(KERNEL_EXTENSIONS):
131+ return False
132+ normalized = os.path.normpath(file_path)
133+ if not normalized.startswith("/"):
134+ normalized = "/" + normalized
135+ return any(seg in normalized for seg in KERNEL_SEGMENTS)
136+ 
137+ 
138+def is_doc_file(file_path: str) -> bool:
139+ """Check whether file_path is an asc-devkit documentation or example."""
140+ if not file_path:
141+ return False
142+ return any(seg in file_path for seg in DOC_SEGMENTS)
143+ 
144+ 
145+def get_covered_categories(docs_read: list) -> set:
146+ """Return the set of foundational category names that have at least one doc read."""
147+ covered = set()
148+ for name, cat in FOUNDATIONAL_CATEGORIES.items():
149+ for doc_path in docs_read:
150+ if any(pattern in doc_path for pattern in cat["patterns"]):
151+ covered.add(name)
152+ break
153+ return covered
154+ 
155+ 
156+def get_missing_categories(docs_read: list) -> dict:
157+ """Return {name: category_dict} for categories that have zero docs read."""
158+ covered = get_covered_categories(docs_read)
159+ return {name: cat for name, cat in FOUNDATIONAL_CATEGORIES.items() if name not in covered}
160+ 
161+ 
162+def emit(response: dict) -> None:
163+ LOGGER.info(json.dumps(response, ensure_ascii=False))
164+ 
165+ 
166+# ── precision gate helpers (unchanged) ───────────────────────
167+def read_precision_state() -> dict:
168+ try:
169+ with open(PRECISION_GATE_FILE) as fh:
170+ return json.load(fh)
171+ except Exception:
172+ return {"d_class_active": False, "edits_allowed": 0}
173+ 
174+ 
175+def write_precision_state(state: dict) -> None:
176+ ensure_state_dir()
177+ with open(PRECISION_GATE_FILE, "w") as fh:
178+ json.dump(state, fh)
179+ 
180+ 
181+def check_precision_gate() -> tuple:
182+ """Check precision gate with D1/D2 stage enforcement.
183+ 
184+ Returns (denied: bool, reason: str | None).
185+ """
186+ state = read_precision_state()
187+ if not state.get("d_class_active", False):
188+ return False, None
189+ 
190+ remaining = state.get("edits_allowed", 0)
191+ if remaining > 0:
192+ state["edits_allowed"] = remaining - 1
193+ write_precision_state(state)
194+ return False, None
195+ 
196+ stage = state.get("stage", "D1")
197+ debug_count = state.get("debug_call_count", 0)
198+ tune_count = state.get("tune_call_count", 0)
199+ total = debug_count + tune_count
200+ 
201+ if total >= 12:
202+ return True, (
203+ "[Precision Gate] TERMINAL: Phase 4 D 类精度修复已耗尽全部 12 次机会\n"
204+ f" D-1 (ascendc-precision-debug): {min(debug_count, 7)}/7 次\n"
205+ f" D-2 (ascendc-precision-tuning): {min(tune_count, 5)}/5 次\n"
206+ "\nPhase 4 失败，请进入 Phase 7 记录 trace。"
207+ )
208+ 
209+ if stage == "D1":
210+ used = debug_count
211+ max_calls = 7
212+ budget = 7
213+ required_skill = "ascendc-precision-debug"
214+ stage_info = f"D-1 阶段: ascendc-precision-debug, 已用 {used}/{max_calls} 次"
215+ if debug_count >= 7:
216+ stage_info += "\n[D-1 已耗尽，下次请调用 ascendc-precision-tuning 进入 D-2 深度审计]"
217+ else:
218+ used = tune_count
219+ max_calls = 5
220+ budget = 5
221+ required_skill = "ascendc-precision-tuning"
222+ stage_info = f"D-2 阶段: ascendc-precision-tuning, 已用 {used}/{max_calls} 次"
223+ 
224+ return True, (
225+ "[Precision Gate] BLOCKED: D 类精度失败禁止直接修改 kernel 代码\n"
226+ "\n"
227+ "evaluate_ascendc.sh 判定当前结果为 D 类（精度不匹配）。\n"
228+ f"{stage_info}\n"
229+ "\n"
230+ f"你必须先调用 {required_skill} Skill 进行结构化根因分析：\n"
231+ " 1. 使用 Skill 工具调用上述 skill\n"
232+ " 2. 等待 Skill 返回诊断结论和修复建议\n"
233+ f" 3. Skill 调用后自动授予 {budget} 次编辑预算用于实施修复\n"
234+ "\n"
235+ f"最近 D 类失败摘要:\n{state.get('failure_summary', '')}"
236+ )
237+ 
238+ 
239+def _deny_doc_gate(coverage, missing, budget):
240+ """Build and emit the doc gate denial message."""
241+ lines = [
242+ "[Doc Gate] BLOCKED: 未完成必备文档查阅就尝试修改 kernel 代码",
243+ "",
244+ f"已覆盖 {coverage}/4 类必备文档。缺少以下类别：",
245+ "",
246+ ]
247+ for i, (name, cat) in enumerate(missing.items(), 1):
248+ lines.append(f" [{i}] {name} — {cat['desc']}")
249+ lines.append(f" 示例: {cat['example']}")
250+ lines.append("")
251+ 
252+ lines.append("请先完成上述文档的查阅，然后再修改 kernel 代码。")
253+ lines.append(f"当前覆盖率: {coverage}/4 → 编辑预算: {budget}")
254+ 
255+ emit({
256+ "hookSpecificOutput": {
257+ "hookEventName": "PreToolUse",
258+ "permissionDecision": "deny",
259+ "permissionDecisionReason": "\n".join(lines),
260+ }
261+ })
262+ 
263+ 
264+def _consume_edit_and_allow(state, coverage, budget, covered):
265+ """Consume one edit from budget and emit allow decision."""
266+ remaining = state.get("edits_remaining", budget)
267+ if coverage == 4 and remaining > budget:
268+ remaining = budget
269+ state["edits_remaining"] = max(remaining - 1, 0)
270+ write_state(state)
271+ 
272+ covered_names = ", ".join(sorted(covered))
273+ emit({
274+ "hookSpecificOutput": {
275+ "hookEventName": "PreToolUse",
276+ "permissionDecision": "allow",
277+ "permissionDecisionReason": (
278+ f"[Doc Gate] 全部 {coverage}/4 类必备文档已覆盖 ({covered_names}), "
279+ f"剩余编辑预算: {state['edits_remaining']}"
280+ ),
281+ }
282+ })
283+ 
284+ 
285+# ── handlers ─────────────────────────────────────────────────
286+def handle_pre(tool_name: str, tool_input: dict) -> None:
287+ """PreToolUse: gate Edit/Write to kernel files.
288+ 
289+ Enforces two independent gates:
290+ 1. Doc Gate — foundational doc categories must be covered
291+ 2. Precision Gate — D-class precision failures require structured debugging
292+ """
293+ if tool_name not in ("Edit", "Write"):
294+ emit({"hookSpecificOutput": {"hookEventName": "PreToolUse", "permissionDecision": "allow"}})
295+ return
296+ 
297+ file_path = tool_input.get("file_path", "")
298+ if not is_kernel_file(file_path):
299+ emit({"hookSpecificOutput": {"hookEventName": "PreToolUse", "permissionDecision": "allow"}})
300+ return
301+ 
302+ state = read_state()
303+ docs_read = state.get("docs_read", [])
304+ covered = get_covered_categories(docs_read)
305+ missing = get_missing_categories(docs_read)
306+ coverage = len(covered)
307+ budget = BUDGET_BY_COVERAGE.get(coverage, 0)
308+ 
309+ if coverage < 4:
310+ _deny_doc_gate(coverage, missing, budget)
311+ return
312+ 
313+ denied, precision_reason = check_precision_gate()
314+ if denied:
315+ emit({
316+ "hookSpecificOutput": {
317+ "hookEventName": "PreToolUse",
318+ "permissionDecision": "deny",
319+ "permissionDecisionReason": precision_reason,
320+ }
321+ })
322+ return
323+ 
324+ _consume_edit_and_allow(state, coverage, budget, covered)
325+ 
326+ 
327+def handle_post(tool_name: str, tool_input: dict) -> None:
328+ """PostToolUse: track doc reads, grant edit budget proportionally."""
329+ if tool_name != "Read":
330+ emit({"hookSpecificOutput": {"hookEventName": "PostToolUse", "permissionDecision": "allow"}})
331+ return
332+ 
333+ file_path = tool_input.get("file_path", "")
334+ if not is_doc_file(file_path):
335+ emit({"hookSpecificOutput": {"hookEventName": "PostToolUse", "permissionDecision": "allow"}})
336+ return
337+ 
338+ state = read_state()
339+ docs_read = state.get("docs_read", [])
340+ 
341+ if file_path not in docs_read:
342+ docs_read.append(file_path)
343+ state["docs_read"] = docs_read
344+ 
345+ # Grant budget based on current coverage
346+ covered = get_covered_categories(docs_read)
347+ coverage = len(covered)
348+ budget = BUDGET_BY_COVERAGE.get(coverage, 0)
349+ missing = get_missing_categories(docs_read)
350+ 
351+ # Reset budget when coverage level changes (don't accumulate)
352+ old_remaining = state.get("edits_remaining", 0)
353+ if coverage == 4:
354+ state["edits_remaining"] = budget
355+ elif old_remaining < budget:
356+ state["edits_remaining"] = budget
357+ 
358+ write_state(state)
359+ 
360+ if coverage < 4:
361+ missing_lines = "\n".join(
362+ f" - {name}: {cat['example']}"
363+ for name, cat in missing.items()
364+ )
365+ msg = (
366+ f"[Doc Gate] 文档已记录 ({coverage}/4 类已覆盖)，编辑预算: {budget}\n"
367+ f"\n还需查阅:\n{missing_lines}"
368+ )
369+ else:
370+ covered_names = ", ".join(sorted(covered))
371+ msg = (
372+ f"[Doc Gate] 全部 {coverage}/4 类必备文档已覆盖 ({covered_names})\n"
373+ f"编辑预算: {budget} 次"
374+ )
375+ 
376+ emit({
377+ "hookSpecificOutput": {
378+ "hookEventName": "PostToolUse",
379+ "permissionDecision": "allow",
380+ "additionalContext": msg,
381+ }
382+ })
383+ 
384+ 
385+def handle_skill_post(tool_name: str, tool_input: dict) -> None:
386+ """PostToolUse for Skill: track precision skill calls and grant edit budget."""
387+ if tool_name != "Skill":
388+ emit({"hookSpecificOutput": {"hookEventName": "PostToolUse", "permissionDecision": "allow"}})
389+ return
390+ 
391+ skill_name = tool_input.get("skill", "")
392+ is_precision = skill_name in ("ascendc-precision-debug", "ascendc-precision-tuning")
393+ if not is_precision:
394+ emit({"hookSpecificOutput": {"hookEventName": "PostToolUse", "permissionDecision": "allow"}})
395+ return
396+ 
397+ state = read_precision_state()
398+ if not state.get("d_class_active", False):
399+ emit({"hookSpecificOutput": {"permissionDecision": "allow"}})
400+ return
401+ 
402+ if skill_name == "ascendc-precision-debug":
403+ state["debug_call_count"] = state.get("debug_call_count", 0) + 1
404+ budget = 7
405+ if state["debug_call_count"] >= 7:
406+ state["stage"] = "D2"
407+ elif skill_name == "ascendc-precision-tuning":
408+ state["tune_call_count"] = state.get("tune_call_count", 0) + 1
409+ budget = 5
410+ state["stage"] = "D2"
411+ 
412+ state["skill_just_called"] = True
413+ state["edits_allowed"] = budget
414+ write_precision_state(state)
415+ 
416+ emit({
417+ "hookSpecificOutput": {
418+ "hookEventName": "PostToolUse",
419+ "permissionDecision": "allow",
420+ "additionalContext": (
421+ f"[Precision Gate] 检测到 {skill_name} Skill 调用完成。\n"
422+ f"已授予 {budget} 次 kernel 编辑预算。请根据 Skill 的修复建议修改代码。"
423+ ),
424+ }
425+ })
426+ 
427+ 
428+# ── main ─────────────────────────────────────────────────────
429+def main() -> None:
430+ args = sys.argv[1:]
431+ mode = "pre"
432+ for a in args:
433+ if a.startswith("--mode="):
434+ mode = a.split("=", 1)[1]
435+ elif a == "--mode" or a == "-m":
436+ idx = args.index(a)
437+ if idx + 1 < len(args):
438+ mode = args[idx + 1]
439+ 
440+ raw = sys.stdin.read()
441+ if not raw or not raw.strip():
442+ emit({"hookSpecificOutput": {"hookEventName": "PreToolUse", "permissionDecision": "allow"}})
443+ return
444+ 
445+ try:
446+ data = json.loads(raw)
447+ except json.JSONDecodeError:
448+ emit({"hookSpecificOutput": {"hookEventName": "PreToolUse", "permissionDecision": "allow"}})
449+ return
450+ 
451+ tool_name = data.get("tool_name", "")
452+ tool_input = data.get("tool_input") or {}
453+ 
454+ if mode == "post":
455+ handle_post(tool_name, tool_input)
456+ elif mode == "skill_post":
457+ handle_skill_post(tool_name, tool_input)
458+ else:
459+ handle_pre(tool_name, tool_input)
460+ 
461+ 
462+if __name__ == "__main__":
463+ main()
Aplugins-community/tilelang2ascendc-ops-generator/hooks/hooks.json+60-0
@@ -0,0 +1,60 @@
1+{
2+ "hooks": {
3+ "SessionStart": [
4+ {
5+ "matcher": "startup|resume|clear|compact",
6+ "hooks": [
7+ {
8+ "type": "command",
9+ "command": "\"${CLAUDE_PLUGIN_ROOT}/hooks/run-hook.cmd\" session-start-tilelang2ascendc-ops-generator",
10+ "async": false
11+ }
12+ ]
13+ }
14+ ],
15+ "PreToolUse": [
16+ {
17+ "matcher": "Bash",
18+ "hooks": [
19+ {
20+ "type": "command",
21+ "command": "python3 \"${CLAUDE_PLUGIN_ROOT}/hooks/skill_script_hook.py\"",
22+ "async": false
23+ }
24+ ]
25+ },
26+ {
27+ "matcher": "Edit|Write",
28+ "hooks": [
29+ {
30+ "type": "command",
31+ "command": "python3 \"${CLAUDE_PLUGIN_ROOT}/hooks/doc_gate.py\" --mode pre",
32+ "async": false
33+ }
34+ ]
35+ }
36+ ],
37+ "PostToolUse": [
38+ {
39+ "matcher": "Read",
40+ "hooks": [
41+ {
42+ "type": "command",
43+ "command": "python3 \"${CLAUDE_PLUGIN_ROOT}/hooks/doc_gate.py\" --mode post",
44+ "async": false
45+ }
46+ ]
47+ },
48+ {
49+ "matcher": "Skill",
50+ "hooks": [
51+ {
52+ "type": "command",
53+ "command": "python3 \"${CLAUDE_PLUGIN_ROOT}/hooks/doc_gate.py\" --mode skill_post",
54+ "async": false
55+ }
56+ ]
57+ }
58+ ]
59+ }
60+}
Aplugins-community/tilelang2ascendc-ops-generator/hooks/run-hook.cmd+17-0
@@ -0,0 +1,17 @@
1+#!/bin/bash
2+set -e
3+ 
4+# Copy run-hook.cmd from ops-direct-invoke (same pattern)
5+# This script is called by hooks.json to dispatch hook commands
6+HOOK_NAME="$1"
7+SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
8+ 
9+case "$HOOK_NAME" in
10+ session-start-tilelang2ascendc-ops-generator)
11+ source "$SCRIPT_DIR/session-start-tilelang2ascendc-ops-generator"
12+ ;;
13+ *)
14+ echo "Unknown hook: $HOOK_NAME"
15+ exit 1
16+ ;;
17+esac
Aplugins-community/tilelang2ascendc-ops-generator/hooks/session-start-tilelang2ascendc-ops-generator+35-0
@@ -0,0 +1,35 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+echo "[tilelang2ascendc-ops-generator] Session started."
12+ 
13+# Precision Gate 中断恢复检查
14+GATE_FILE="${CLAUDE_PLUGIN_ROOT}/.claude/hooks/state/precision_gate.json"
15+if [ -f "$GATE_FILE" ]; then
16+ D_CLASS_ACTIVE=$(python3 -c "import json; d=json.load(open('$GATE_FILE')); print(d.get('d_class_active', False))" 2>/dev/null)
17+ if [ "$D_CLASS_ACTIVE" = "True" ]; then
18+ OUTPUT_DIR=$(python3 -c "import json; d=json.load(open('$GATE_FILE')); print(d.get('output_dir', 'unknown'))" 2>/dev/null)
19+ STAGE=$(python3 -c "import json; d=json.load(open('$GATE_FILE')); print(d.get('stage', 'D1'))" 2>/dev/null)
20+ DEBUG_COUNT=$(python3 -c "import json; d=json.load(open('$GATE_FILE')); print(d.get('debug_call_count', 0))" 2>/dev/null)
21+ TUNE_COUNT=$(python3 -c "import json; d=json.load(open('$GATE_FILE')); print(d.get('tune_call_count', 0))" 2>/dev/null)
22+ TOTAL=$((DEBUG_COUNT + TUNE_COUNT))
23+ echo ""
24+ echo "╔══════════════════════════════════════════════════════════════╗"
25+ echo "║ [Precision Gate] D 类精度失败待处理 ║"
26+ echo "╠══════════════════════════════════════════════════════════════╣"
27+ echo "║ output_dir : $OUTPUT_DIR"
28+ echo "║ stage : $STAGE (debug=$DEBUG_COUNT, tune=$TUNE_COUNT, total=$TOTAL/12)"
29+ echo "║ 状态 : 上一轮 Subagent 返回 D 类精度不匹配 ║"
30+ echo "║ Orchestrator 必须: 立即重新调用 ascend-kernel-developer ║"
31+ echo "║ Subagent 进入 D 类精度修复流程 ║"
32+ echo "╚══════════════════════════════════════════════════════════════╝"
33+ echo ""
34+ fi
35+fi
Aplugins-community/tilelang2ascendc-ops-generator/hooks/skill_script_hook.py+362-0
@@ -0,0 +1,362 @@
1+#!/usr/bin/env python3
2+# ----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------------------------------------
11+"""
12+Skill 脚本执行 Hook (PreToolUse 协议)。
13+ 
14+对 Bash 工具调用进行拦截：
15+- 命中 INTERCEPTED_PATTERNS: 直接由本 hook 执行脚本，并以 PreToolUse 协议返回
16+ permissionDecision=deny + additionalContext，把执行结果（exit_code/stdout/stderr）
17+ 注入回 agent 的上下文，阻止 harness 二次执行。
18+- 不命中: 输出 permissionDecision=allow，让 harness 正常执行（不在 hook 中重复执行）。
19+ 
20+输入：从 stdin 读取 harness PreToolUse JSON：
21+ {"tool_name": "Bash", "tool_input": {"command": "..."}, ...}
22+ 
23+输出：单行 PreToolUse JSON 到 stdout。
24+"""
25+ 
26+import json
27+import logging
28+import os
29+import re
30+import subprocess
31+import sys
32+import time
33+from pathlib import Path
34+ 
35+ 
36+_protocol_logger = logging.getLogger("skill_script_hook.protocol")
37+_protocol_logger.propagate = False
38+_protocol_logger.setLevel(logging.INFO)
39+_protocol_handler = logging.StreamHandler(sys.stdout)
40+_protocol_handler.setFormatter(logging.Formatter("%(message)s"))
41+_protocol_logger.addHandler(_protocol_handler)
42+ 
43+ 
44+# 每个条目: (interpreter_regex, script_basename_regex)
45+# 只匹配“某个解释器 + 脚本路径”这种真正的脚本调用，
46+# 不匹配命令字符串里仅出现脚本名的情况（如 cat / grep / echo JSON 等）。
47+_PY_INTERP = r"(?:python|python3|python3\.\d+)"
48+_SH_INTERP = r"(?:bash|sh)"
49+ 
50+INTERCEPTED_PATTERNS = [
51+ (_PY_INTERP, r"validate_tilelang_impl\.py"),
52+ (_PY_INTERP, r"validate_ascendc_impl\.py"),
53+ (_SH_INTERP, r"evaluate_tilelang\.sh"),
54+ (_SH_INTERP, r"evaluate_ascendc\.sh"),
55+ (_PY_INTERP, r"performance\.py"),
56+ (_PY_INTERP, r"build_ascendc\.py"),
57+ (_PY_INTERP, r"verification_ascendc\.py"),
58+ (_PY_INTERP, r"verification_tilelang\.py"),
59+]
60+ 
61+# 命令前缀允许的部分（环境变量赋值、cd ... &&、source ... &&）
62+# 这里只是用于推断"真实的脚本调用"，不需要覆盖所有 shell 语法。
63+_LEADING_PREFIX = (
64+ r"^\s*"
65+ r"(?:export\s+)?" # 可选的 export 关键字
66+ r"(?:[A-Za-z_][A-Za-z0-9_]*=\S+\s+)*" # ENV=VAL ...
67+ r"(?:cd\s+\S+\s*&&\s*)?" # cd <dir> &&
68+ r"(?:export\s+)?" # 可选的 export 关键字
69+ r"(?:[A-Za-z_][A-Za-z0-9_]*=\S+\s+)*" # 再次允许 ENV=VAL
70+ r"(?:&&\s+)*" # 支持 && 链式连接
71+)
72+ 
73+# 项目根目录：优先从环境变量获取，fallback 到 __file__ 计算
74+# 避免 hook 被以相对路径调用时 __file__ 解析错误
75+PROJECT_ROOT = os.environ.get("PROJECT_ROOT")
76+if not PROJECT_ROOT:
77+ PROJECT_ROOT = os.path.dirname(os.path.dirname(os.path.dirname(os.path.abspath(__file__))))
78+ 
79+ 
80+def should_intercept(command: str) -> bool:
81+ # 仅当命令的"实际执行段"是 <interpreter> <path-to-script> 形式时才拦截。
82+ # 兼容：环境变量前缀、可选的 `cd <dir> &&` 前缀。
83+ for interp, script in INTERCEPTED_PATTERNS:
84+ # 路径里允许出现 / 和非空白字符；脚本名用 \b 边界
85+ pattern = (
86+ _LEADING_PREFIX
87+ + interp
88+ + r"\s+"
89+ + r"(?:\S*/)?"
90+ + script
91+ + r"(?:\s|$)"
92+ )
93+ if re.search(pattern, command):
94+ return True
95+ return False
96+ 
97+ 
98+def emit_pretooluse(permission_decision: str, *, reason: str = "", additional_context: str = ""):
99+ output = {
100+ "hookSpecificOutput": {
101+ "hookEventName": "PreToolUse",
102+ "permissionDecision": permission_decision,
103+ }
104+ }
105+ if reason:
106+ output["hookSpecificOutput"]["permissionDecisionReason"] = reason
107+ if additional_context:
108+ output["hookSpecificOutput"]["additionalContext"] = additional_context
109+ _protocol_logger.info(json.dumps(output, ensure_ascii=False))
110+ 
111+ 
112+def truncate(text: str, limit: int = 8000) -> str:
113+ if not text:
114+ return ""
115+ if len(text) <= limit:
116+ return text
117+ head = text[: limit // 2]
118+ tail = text[-limit // 2:]
119+ return f"{head}\n\n... [truncated {len(text) - limit} chars] ...\n\n{tail}"
120+ 
121+ 
122+def _run_command(command: str, cwd: str):
123+ """Execute a shell command and return (exit_code, stdout, stderr, duration_ms)."""
124+ start_time = time.time()
125+ try:
126+ import shlex
127+ cmd_parts = shlex.split(command)
128+ proc = subprocess.run(
129+ cmd_parts, shell=False, cwd=cwd,
130+ capture_output=True, text=True, timeout=1800,
131+ )
132+ exit_code = proc.returncode
133+ stdout = proc.stdout or ""
134+ stderr = proc.stderr or ""
135+ except subprocess.TimeoutExpired as e:
136+ exit_code = 124
137+ stdout = (e.stdout if e.stdout else "")
138+ stderr = (e.stderr if e.stderr else "")
139+ stderr += "\n[HOOK] 命令执行超时（30 分钟）"
140+ except Exception as e:
141+ exit_code = 1
142+ stdout = ""
143+ stderr = f"[HOOK] 执行异常: {type(e).__name__}: {e}"
144+ 
145+ duration_ms = int((time.time() - start_time) * 1000)
146+ return exit_code, stdout, stderr, duration_ms
147+ 
148+# ── Precision Gate (D-class enforcement) ────────────────────
149+_PRECISION_GATE_STATE = os.path.join(
150+ os.path.dirname(os.path.dirname(os.path.abspath(__file__))),
151+ ".claude", "hooks", "state", "precision_gate.json",
152+)
153+ 
154+# Patterns for D-class detection in evaluate_ascendc.sh output
155+_A_CLASS_ERROR_PATTERNS = [
156+ r"\berror:\s", # compilation error
157+ r"\bfatal error:\s", # fatal compilation error
158+ r"undefined reference", # linker error
159+ r"Segmentation fault",
160+ r"core dumped",
161+ r"cannot find -l", # library not found
162+ r"No such file or directory",
163+ r"CMake Error",
164+ r"make\[\d+\]: \*\*\*", # make build error
165+]
166+_A_CLASS_PATTERN_RE = re.compile("|".join(_A_CLASS_ERROR_PATTERNS), re.IGNORECASE)
167+ 
168+_D_CLASS_PASS_RE = re.compile(r"^\s*Result:\s*(pass|fail)\s*$", re.MULTILINE | re.IGNORECASE)
169+_D_CLASS_NUMERIC_RE = re.compile(r"max_abs_diff\s*=\s*[\d.e+\-]+|MERE\s*=\s*[\d.e+\-]+|matched_ratio\s*=\s*[\d.]+")
170+ 
171+_EVAL_ASCENDC_SCRIPT_RE = re.compile(r"evaluate_ascendc\.sh")
172+ 
173+ 
174+def _is_evaluate_ascendc(command: str) -> bool:
175+ return bool(_EVAL_ASCENDC_SCRIPT_RE.search(command))
176+ 
177+ 
178+def _extract_output_dir(command: str) -> str:
179+ """Extract output_dir argument from evaluate_ascendc.sh command."""
180+ m = re.search(r"evaluate_ascendc\.sh\s+(\S+)", command)
181+ return m.group(1) if m else ""
182+ 
183+ 
184+def _classify_result(exit_code: int, stdout: str, stderr: str) -> str:
185+ """
186+ Classify evaluate_ascendc.sh result as "D" (precision failure),
187+ "A" (code/environment error), or "PASS".
188+ """
189+ combined = f"{stdout}\n{stderr}"
190+ 
191+ # Check for PASS
192+ pass_matches = _D_CLASS_PASS_RE.findall(combined)
193+ if pass_matches:
194+ last_result = pass_matches[-1].lower()
195+ if last_result == "pass":
196+ return "PASS"
197+ 
198+ # Check for A-class (compilation/runtime error)
199+ if _A_CLASS_PATTERN_RE.search(combined):
200+ return "A"
201+ 
202+ # Check for D-class: has numeric diff output and Result: fail
203+ if _D_CLASS_NUMERIC_RE.search(combined):
204+ return "D"
205+ 
206+ # Non-zero exit without clear A-class or D-class patterns → A
207+ if exit_code != 0:
208+ return "A"
209+ 
210+ return "UNKNOWN"
211+ 
212+ 
213+def _ensure_precision_gate_dir() -> None:
214+ os.makedirs(os.path.dirname(_PRECISION_GATE_STATE), exist_ok=True)
215+ 
216+ 
217+def _read_precision_gate() -> dict:
218+ """Read existing precision gate state to preserve counters across iterations."""
219+ try:
220+ with open(_PRECISION_GATE_STATE) as fh:
221+ return json.load(fh)
222+ except Exception:
223+ return {}
224+ 
225+ 
226+def _update_precision_gate(classification: str, stdout: str, output_dir: str) -> None:
227+ """Write or clear precision gate state based on evaluate_ascendc.sh result.
228+ On D-class: resets edits_allowed/skill_just_called but PRESERVES
229+ stage/debug_call_count/tune_call_count across iterations.
230+ On PASS/A: full reset.
231+ """
232+ if classification == "D":
233+ # Extract failure summary (first few lines of case diffs)
234+ summary_lines = []
235+ for line in stdout.splitlines():
236+ if "max_abs_diff" in line or "Result:" in line:
237+ summary_lines.append(line.strip())
238+ failure_summary = "\n".join(summary_lines[:5])
239+ 
240+ # Preserve stage and counters from previous iteration
241+ existing = _read_precision_gate()
242+ stage = existing.get("stage", "D1")
243+ debug_count = existing.get("debug_call_count", 0)
244+ tune_count = existing.get("tune_call_count", 0)
245+ 
246+ # Auto-advance stage based on accumulated counts
247+ if debug_count >= 7:
248+ stage = "D2"
249+ 
250+ state = {
251+ "d_class_active": True,
252+ "output_dir": output_dir,
253+ "edits_allowed": 0,
254+ "skill_just_called": False,
255+ "stage": stage,
256+ "debug_call_count": debug_count,
257+ "tune_call_count": tune_count,
258+ "failure_summary": failure_summary,
259+ }
260+ else:
261+ # PASS or A-class: full reset
262+ state = {"d_class_active": False, "edits_allowed": 0}
263+ 
264+ _ensure_precision_gate_dir()
265+ with open(_PRECISION_GATE_STATE, "w") as fh:
266+ json.dump(state, fh)
267+ 
268+ 
269+def execute_intercepted(command: str) -> None:
270+ cwd = os.getcwd()
271+ if os.path.isdir(PROJECT_ROOT):
272+ cwd = PROJECT_ROOT
273+ 
274+ exit_code, stdout, stderr, duration_ms = _run_command(command, cwd)
275+ 
276+ # Update precision gate state if this is an evaluate_ascendc.sh call
277+ is_eval_ascendc = _is_evaluate_ascendc(command)
278+ classification = None
279+ if is_eval_ascendc:
280+ classification = _classify_result(exit_code, stdout, stderr)
281+ output_dir = _extract_output_dir(command)
282+ _update_precision_gate(classification, stdout, output_dir)
283+ 
284+ additional_context = (
285+ f"[skill_script_hook intercepted execution]\n"
286+ f"command: {command}\n"
287+ f"cwd: {cwd}\n"
288+ f"exit_code: {exit_code}\n"
289+ f"duration_ms: {duration_ms}\n"
290+ f"--- stdout ---\n{truncate(stdout)}\n"
291+ f"--- stderr ---\n{truncate(stderr)}\n"
292+ )
293+ 
294+ status = "成功" if exit_code == 0 else "失败"
295+ 
296+ class_note = ""
297+ if is_eval_ascendc and classification:
298+ class_labels = {"D": "D类-精度不匹配", "A": "A类-代码/编译错误", "PASS": "通过"}
299+ label = class_labels.get(classification, classification)
300+ class_note = f" | 错误分类: {label}"
301+ if classification == "D":
302+ class_note += (
303+ " | [precision_gate] kernel 编辑已锁定，修改代码前必须先调用 "
304+ "ascendc-precision-debug 或 ascendc-precision-tuning Skill"
305+ )
306+ 
307+ reason = (
308+ f"[Hook 拦截提示] 该命令命中 skill_script_hook 拦截规则，"
309+ f"已由 hook 代为执行（{status}，exit_code={exit_code}）。{class_note}"
310+ f"执行结果见下方 additionalContext，原命令已被替换为 no-op，不会重复执行。"
311+ )
312+ 
313+ output = {
314+ "hookSpecificOutput": {
315+ "hookEventName": "PreToolUse",
316+ "permissionDecision": "allow",
317+ "permissionDecisionReason": reason,
318+ "additionalContext": additional_context,
319+ "updatedInput": {
320+ "command": f"echo '[hook-noop] original command was intercepted and executed by skill_script_hook'"
321+ },
322+ }
323+ }
324+ _protocol_logger.info(json.dumps(output, ensure_ascii=False))
325+ 
326+ 
327+def read_stdin_command():
328+ try:
329+ raw = sys.stdin.read()
330+ except Exception:
331+ return None, None
332+ if not raw or not raw.strip():
333+ return None, None
334+ try:
335+ data = json.loads(raw)
336+ except json.JSONDecodeError:
337+ return None, None
338+ if not isinstance(data, dict):
339+ return None, None
340+ tool_name = data.get("tool_name")
341+ tool_input = data.get("tool_input") or {}
342+ command = tool_input.get("command") if isinstance(tool_input, dict) else None
343+ return tool_name, command
344+ 
345+ 
346+def main():
347+ tool_name, command = read_stdin_command()
348+ 
349+ # 没有命令或不是 Bash —— allow，由 harness 处理
350+ if not command or tool_name != "Bash":
351+ emit_pretooluse("allow")
352+ return
353+ 
354+ if should_intercept(command):
355+ execute_intercepted(command)
356+ else:
357+ # 不拦截 —— 让 harness 正常处理（不在此重复执行）
358+ emit_pretooluse("allow")
359+ 
360+ 
361+if __name__ == "__main__":
362+ main()
Aplugins-community/tilelang2ascendc-ops-generator/quickstart.md+242-0
@@ -0,0 +1,242 @@
1+# CANNBot Ascend C 算子自动生成快速入门指南
2+ 
3+## 概述
4+ 
5+Ascend C 算子自动生成模式适用于**从 PyTorch Model 自动生成 Ascend C Kernel**场景，支持双路径：简单算子走 ops-direct-invoke 工作流（Architect 设计 → Developer 实现 → Reviewer 审查），复杂算子走 TileLang 设计表达 → AscendC 转译优化，端到端完成算子开发。
6+ 
7+### 工作流
8+ 
9+```
10+Phase 0: 参数确认 + 算子分类 (解析输入，判定简单/复杂路径)
11+Phase 1: 环境准备 + 工程初始化 (复制算子文件 + 初始化 kernel 工程)
12+Phase 2: Case 精简 (tilelang2ascend-case-simplifier)
13+Phase 3: 设计表达 (分支)
14+ ├─ 简单算子: 架构设计 + 设计串讲 (ops-direct-invoke: DESIGN.md + PLAN.md + WALKTHROUGH.md)
15+ └─ 复杂算子: TileLang 设计 (tilelang2ascend-tilelang-designer + 退化检测 + 迭代)
16+Phase 4: AscendC 生成与验证 (分支)
17+ ├─ 简单算子: 开发实现 + 代码审查 + 修复循环 (ops-direct-invoke: 渐进式开发 + REVIEW.md + 最多3轮修复)
18+ └─ 复杂算子: 转译 (tilelang2ascend-translator + 退化检测 + 迭代)
19+Phase 5: 性能分析 (ops-profiling --compare 模式)
20+Phase 6: 全量验证
21+Phase 7: Trace 记录 + 知识演进 (tilelang2ascend-trace-recorder: trace.md →
22+ 经 ops-knowledge-ingest 路由按 OKF 格式写入共享知识库
23+ runbooks/，供后续任务经 knowledge-query 检索复用)
24+```
25+ 
26+### 算子分类路由
27+ 
28+| 路径 | 算子类型 | Skill 链 |
29+|------|---------|---------|
30+| 简单算子 | Index, IndexPut, Gather, Scatter, Nonzero, RepeatInterleave, EmbeddingDenseBackward | ops-direct-invoke（Architect → Developer → Reviewer） |
31+| 复杂算子 | Attention, MatMul 变体, Norm 变体, Sort, TopK, 多输入融合 | tilelang2ascend-tilelang-designer → tilelang2ascend-translator |
32+ 
33+## 一、环境搭建
34+ 
35+### 前置条件
36+ 
37+- 已安装 Python 3.8 或更高版本
38+- 已安装 PyTorch 2.0 或更高版本
39+- 已安装 tilelang-ascend，具体安装方式请查阅 [tilelang-ascend readme](https://github.com/tile-ai/tilelang-ascend/blob/ascendc_pto/README.md#method-3-compile-and-install-from-source)
40+- 已安装 CANN Toolkit（建议 ≥ 9.0.0），具体版本配套关系请查阅 [CANN Release Notes](https://www.hiascend.com/cann/document)
41+- 已配置 NPU 设备（支持 Ascend 910/950 PR 等芯片）
42+- 已安装 OpenCode、Claude Code、TRAE、Cursor、Copilot、CodeArts 等受支持的 AI 编程工具
43+ 
44+### Claude Code（推荐）
45+ 
46+**首选：Plugin Marketplace（一键安装）**
47+ 
48+```bash
49+# 注册 marketplace（首次，GitCode 仓库需完整 URL）
50+/plugin marketplace add https://gitcode.com/cann/cannbot-skills.git
51+ 
52+# 安装插件
53+/plugin install tilelang2ascendc-ops-generator@cannbot
54+```
55+ 
56+**备选：init.sh 脚本**
57+ 
58+```bash
59+git clone https://gitcode.com/cann/cannbot-skills.git
60+cd cannbot-skills/plugins-community/tilelang2ascendc-ops-generator
61+bash init.sh project claude # 项目级
62+bash init.sh global claude # 全局级
63+```
64+ 
65+### 其他工具
66+ 
67+<details>
68+<summary> OpenCode</summary>
69+ 
70+```bash
71+git clone https://gitcode.com/cann/cannbot-skills.git
72+cd cannbot-skills/plugins-community/tilelang2ascendc-ops-generator
73+bash init.sh project opencode # 项目级（默认）
74+bash init.sh global opencode # 全局级
75+```
76+ 
77+</details>
78+ 
79+<details>
80+<summary>TRAE</summary>
81+ 
82+仅支持项目级安装。
83+ 
84+```bash
85+git clone https://gitcode.com/cann/cannbot-skills.git
86+cd cannbot-skills/plugins-community/tilelang2ascendc-ops-generator
87+bash init.sh project trae
88+```
89+ 
90+</details>
91+ 
92+<details>
93+<summary>Cursor</summary>
94+ 
95+```bash
96+git clone https://gitcode.com/cann/cannbot-skills.git
97+cd cannbot-skills/plugins-community/tilelang2ascendc-ops-generator
98+bash init.sh project cursor # 项目级
99+bash init.sh global cursor # 全局级
100+```
101+ 
102+安装后在项目根目录生成 `.cursor/` 目录（项目级）或 `~/.cursor/` 目录（全局级），包含 skills/、agents/ 和 AGENTS.md。
103+ 
104+</details>
105+ 
106+<details>
107+<summary>Copilot</summary>
108+ 
109+```bash
110+git clone https://gitcode.com/cann/cannbot-skills.git
111+cd cannbot-skills/plugins-community/tilelang2ascendc-ops-generator
112+bash init.sh project copilot # 项目级
113+bash init.sh global copilot # 全局级
114+```
115+ 
116+安装后在项目根目录生成 `.github/` 目录（项目级）或 `~/.copilot/` 目录（全局级），包含 skills/、agents/ 和 AGENTS.md。
117+ 
118+</details>
119+ 
120+<details>
121+<summary>CodeArts</summary>
122+ 
123+```bash
124+git clone https://gitcode.com/cann/cannbot-skills.git
125+cd cannbot-skills/plugins-community/tilelang2ascendc-ops-generator
126+bash init.sh project codearts # 项目级
127+bash init.sh global codearts # 全局级
128+```
129+ 
130+安装后在项目根目录生成 `.codeartsdoer/` 目录（项目级）或 `~/.codeartsdoer/` 目录（全局级），包含 skills/、agents/ 和 AGENTS.md。
131+ 
132+</details>
133+ 
134+## 二、快速上手
135+ 
136+### 使用方式
137+ 
138+#### 场景一：单算子生成
139+在交互界面中输入算子开发需求：
140+ 
141+```
142+生成ascendC算子，npu=0，算子描述文件为 /path/to/op/model.py，输出到 /path/to/output/op_name/
143+```
144+ 
145+CANNBot 会自动调度 ascend-kernel-developer 按 7 Phase 流程执行：
146+1. 解析参数、判定算子类型（简单/复杂）
147+2. 准备环境、初始化 kernel 工程
148+3. 精简测试用例（≤ 10 个代表性 case）
149+4. 设计表达（简单算子走 ops-direct-invoke 架构设计 + 设计串讲，复杂算子做 TileLang 设计）
150+5. AscendC 生成与验证（简单算子走开发实现 + 代码审查 + 修复循环，复杂算子走转译 + 迭代验证）
151+6. 性能分析（对比 reference 和 ascendc）
152+7. 全量用例验证 + 生成 trace.md 并完成知识演进（共享知识库 runbooks/ 更新）
153+ 
154+#### 场景二：批量性能测试
155+ 
156+适用于对多个已生成算子进行批量性能对比测试，支持单 NPU 串行或多 NPU 并行执行。
157+ 
158+**操作步骤**：
159+ 
160+1. 确保每个算子目录下已有 `model.py`、`model_new_ascendc.py` 和测试用例文件
161+2. 执行批量测试脚本：
162+ 
163+```bash
164+bash .claude/skills/ops-profiling/scripts/msprof_profile_run.sh --batch \
165+ --base-dir /path/to/output_performance \
166+ --max-jobs 7 \
167+ --device-start 1
168+```
169+ 
170+**参数说明**：
171+- `--base-dir`: 包含多个算子输出子目录的根目录（必填）
172+- `--max-jobs`: 最大并发数（默认 7）
173+- `--device-start`: 起始 NPU 设备 ID（默认 1）
174+- `--warm-up`: 预热次数（默认 3）
175+ 
176+### 产出物示例
177+ 
178+```
179+{output_dir}/
180+├── model.py # 算子描述文件（只读）
181+├── <op_name>.json # 原始测试用例文件（备份保留）
182+├── <op_name>.json.bak # 原始 .json 备份
183+├── docs/
184+│ ├── DESIGN.md # 架构设计文档（简单算子路径）
185+│ ├── PLAN.md # 开发计划（简单算子路径）
186+│ ├── WALKTHROUGH.md # 设计串讲记录（简单算子路径）
187+│ └── REVIEW.md # 代码审查报告（简单算子路径）
188+├── design/
189+│ ├── block_level/ # Block-level 设计（复杂算子路径）
190+│ └── tile_level/ # Tile-level 设计（复杂算子路径）
191+├── kernel/ # AscendC kernel 实现
192+├── model_new_tilelang.py # TileLang 实现（复杂算子路径）
193+├── model_new_ascendc.py # AscendC 实现
194+├── performance.json # 性能数据
195+└── trace.md # 执行记录
196+```
197+ 
198+## 三、可用技能
199+ 
200+| Skill | 用途 | 适用路径 |
201+|-------|------|---------|
202+| `tilelang2ascend-case-simplifier` | 测试用例精简 | 所有 |
203+| ops-direct-invoke (Architect + Developer + Reviewer) | 简单算子架构设计、开发实现、代码审查 | 简单算子 |
204+| `tilelang2ascend-tilelang-designer` | TileLang kernel 设计表达 | 复杂算子 |
205+| `tilelang2ascend-translator` | TileLang→AscendC 转译 | 复杂算子 |
206+| `tilelang2ascend-operator-project-init` | Kernel 工程初始化 | 所有 |
207+| `ops-profiling` | 性能测试分析（--compare 模式） | 所有 |
208+| `tilelang2ascend-trace-recorder` | 执行记录生成 | 所有 |
209+ 
210+## 四、常见问题
211+ 
212+### Q: 如何查看帮助信息？
213+ 
214+```bash
215+bash init.sh --help
216+```
217+ 
218+### Q: 项目级和全局安装如何选择？
219+ 
220+- **项目级**：适合多项目开发，每个项目可以有不同配置
221+- **全局**：适合单一项目，全局生效
222+ 
223+### Q: 简单算子和复杂算子有什么区别？
224+ 
225+简单算子（Index, Gather, Scatter 等）走 ops-direct-invoke 工作流：Architect 架构设计 → Developer 渐进式开发 → Reviewer 代码审查（100分制评分 + 修复循环）。复杂算子（Attention, MatMul, Norm, Sort 等）需要先做 TileLang 设计表达，再转译为 AscendC 代码。
226+ 
227+### Q: TileLang 验证失败了怎么办？
228+ 
229+TileLang 当前主要用于设计表达，不是 correctness gate。若 TileLang 验证失败但设计意图正确，可跳过 TileLang 验证直接进入 AscendC 转译阶段。在 trace.md 中会记录跳过原因。
230+ 
231+### Q: AscendC 验证迭代上限是多少？
232+ 
233+最多 3 轮迭代。超过上限后停止并报告当前状态。
234+ 
235+---
236+ 
237+## 总结
238+ 
239+1. 从 PyTorch Model 出发，支持双路径（简单/复杂）端到端完成算子开发
240+2. 使用 `init.sh` 脚本一键安装（OpenCode 推荐），Claude Code 用户也可用 `/plugin install` 一键安装
241+3. 7 Phase 工作流（含 Phase 7 知识演进闭环），自动算子分类路由，退化检测，迭代修复
242+4. 产出物包含设计文档/设计表达、kernel 代码、性能报告、trace 记录和演进知识库条目
Aplugins-community/tilelang2ascendc-ops-generator/workflows/development-guide.md+161-0
@@ -0,0 +1,161 @@
1+# AscendC Ops Developer 开发指南
2+ 
3+## 概述
4+ 
5+本 Team 实现从 PyTorch Model 到 AscendC Kernel 的端到端自动开发流程。
6+ 
7+## 工作流架构
8+ 
9+```
10+用户输入 (npu, op_file, output_dir)
11+ │
12+ ▼
13+CANNBot Primary Agent (AGENTS.md)
14+ │ 解析参数、调度 Subagent
15+ ▼
16+ascend-kernel-developer Subagent
17+ │
18+ ├── Phase 0: 参数确认 + 算子分类（简单/复杂）
19+ ├── Phase 1: 环境准备
20+ ├── Phase 2: Case 精简 (tilelang2ascend-case-simplifier)
21+ ├── Phase 3: 设计表达（分支）
22+ │ ├─ 简单: ops-direct-invoke 架构设计 + 设计串讲
23+ │ └─ 复杂: TileLang 设计 (tilelang2ascend-tilelang-designer) ← 迭代循环
24+ ├── Phase 4: AscendC 生成与验证（分支）
25+ │ ├─ 简单: ops-direct-invoke 开发实现 + 代码审查 + 修复循环（3轮上限）
26+ │ └─ 复杂: AscendC 转译 (tilelang2ascend-translator) ← 迭代循环（3轮上限）
27+ ├── Phase 5: 性能分析 (ops-profiling --compare 模式)
28+ ├── Phase 6: 全量验证
29+ └── Phase 7: Trace 记录 + 知识演进 (tilelang2ascend-trace-recorder:
30+ trace.md → 经 ops-knowledge-ingest 路由写入共享知识库 runbooks/)
31+```
32+ 
33+## Phase 3-4 迭代机制
34+ 
35+### Phase 3: TileLang 设计迭代
36+ 
37+```
38+generation → AST退化检测 → [通过] → 功能验证 → [通过] → 完成
39+ ↓ 失败 ↓ 失败
40+ Conductor分析 → 生成修复建议 → 重新生成
41+ ↑
42+ 最多 5 轮迭代 ─┘
43+```
44+ 
45+退化子类型：
46+- Type1: 无 TileLang kernel 导入（纯 PyTorch）
47+- Type2: 有导入但 forward() 未调用
48+- Type3: 调用 kernel 但部分计算仍用 PyTorch
49+- Type4: forward() 中存在逐元素 for 循环
50+ 
51+### Phase 4: AscendC 转译迭代
52+ 
53+```
54+generation → AST退化检测 → [通过] → 功能验证 → [通过] → 完成
55+ ↓ 失败 ↓ 失败
56+ Conductor分析 → 生成修复建议 → 重新生成
57+ ↑
58+ 最多 3 轮迭代 ─┘
59+```
60+ 
61+退化子类型：
62+- Type1: 无 AscendC 扩展导入（纯 PyTorch）
63+- Type2: 有导入但 forward() 未调用 kernel
64+- Type3: 调用 kernel 但部分计算仍用 PyTorch
65+- Type4: forward() 中存在逐元素 for 循环
66+ 
67+## 错误分类
68+ 
69+| 分类 | 含义 | 处理 |
70+|------|------|------|
71+| A 类 | 代码逻辑/算法错误（可修复） | 生成修复建议，继续迭代 |
72+| B 类 | 环境/基础设施错误（不可修复） | 立即终止 |
73+| C 类 | 同一 A 类子类型连续 ≥ 3 次 | 立即终止 |
74+ 
75+## 产出物
76+ 
77+```
78+{output_dir}/
79+├── model.py # 算子描述（只读）
80+├── <op_name>.json # 原始测试用例文件（备份保留）
81+├── <op_name>.json.bak # 原始 .json 备份
82+├── design/
83+│ ├── block_level/ # Block-level 设计
84+│ └── tile_level/ # Tile-level 设计
85+├── kernel/ # AscendC kernel
86+├── model_new_tilelang.py # TileLang 实现
87+├── model_new_ascendc.py # AscendC 实现
88+├── preformance.json # 性能数据
89+└── trace.md # 执行记录（Phase 7 知识演进输入）
90+```
91+ 
92+## 知识演进闭环
93+ 
94+- 每个任务完成后，`tilelang2ascend-trace-recorder` 在 Phase 7 从 trace.md 的走偏点/
95+ 行为轮次/关键错误中提取可复用经验，经 `ops-knowledge-ingest` 标准路由按
96+ OKF (okf.v1) 格式写入共享知识库（`$CANNBOT_KNOWLEDGE_ROOT`，当前
97+ `/home/asc-gen-knowledge`）的 `runbooks/` 树（field_notes/optimization 子树；
98+ 维度分类 sync/api/tiling/precision/process 走 tags，不建目录），
99+ 并同步三件套（逐层 index.md + 检索索引/图谱重建 + log）。
100+- 后续任务的 Phase 0 与 Phase 4 转译前经 `knowledge-query` skill 检索知识库
101+ （preflight/search，禁自维护索引），命中"触发条件"即规避已知坑，
102+ 从而降低迭代轮数、提升生成质量。
103+- 知识库卡片可随平台升级失效；失效时以官方文档为准并标记 superseded。
104+ 
105+## 关键约束
106+ 
107+- TileLang 仅用于设计表达，不作为 correctness/performance gate
108+- model_new_*.py 中禁止使用 torch 算子
109+- 必须将核心计算融合成单个算子
110+- 文件操作范围限制在 {output_dir}/ 内
111+- 优先使用块级/向量化操作，避免标量逐元素写法
112+ 
113+## 评测环境清理最佳实践
114+ 
115+### 问题
116+ 
117+`npu-kernelbench` 的 `pre_build()` 使用 `mkdir(exist_ok=True)`，不清理旧工作目录。连续运行时残留三类有害物：
118+ 
119+| 残留物 | 路径 | 影响 |
120+|--------|------|------|
121+| Profile 数据累积 | `profile/<op>-<sol>/<pid>_<timestamp>_ascend_pt/` | 每次运行追加新子目录，profiler 可能读到旧数据，导致延迟测量错误（出现固定异常值） |
122+| CMake 构建缓存 | `build/CMakeCache.txt`, `build/CMakeFiles/` | 可能跳过重新配置，复用旧编译参数 |
123+| traces.jsonl 追加模式 | 默认追加 | 旧 trace 混入新结果 |
124+ 
125+### 方案：运行前清理（推荐）
126+ 
127+评测脚本（`evaluate_ascendc.sh` / `evaluate_tilelang.sh`）执行前，必须清理工作目录和旧 traces：
128+ 
129+```bash
130+# 清理工作目录（CMake 缓存 + 旧 kernel.so + profile 数据）
131+WORK_DIR="/tmp/npu-kernelbench/${OP_NAME}-${SOLUTION_NAME}"
132+rm -rf "$WORK_DIR"
133+ 
134+# 清理旧 traces（避免追加模式混入旧结果）
135+TRACE_FILE="/tmp/npu-kernelbench/traces.jsonl"
136+rm -f "$TRACE_FILE"
137+```
138+ 
139+也可封装为一键脚本：
140+ 
141+```bash
142+#!/bin/bash
143+set -euo pipefail
144+# 1. 清理工作目录
145+rm -rf /tmp/npu-kernelbench/${OP_NAME}-${SOLUTION_NAME}
146+# 2. 清理旧 traces
147+rm -f /tmp/npu-kernelbench/traces.jsonl
148+# 3. source CANN 环境 + 运行评测
149+source <CANN_PATH>/set_env.sh
150+cd <BENCH_DIR>
151+uv run npu-kernelbench eval \
152+ --definition "$DEF" --workload "$WL" --solution "$SOL" \
153+ -o /tmp/npu-kernelbench/traces.jsonl
154+```
155+ 
156+### 验证清理效果
157+ 
158+清理后运行评测，确认：
159+- `traces.jsonl` 的行数 = workload 的 case 数（无多余旧记录）
160+- profile 目录下只有本次运行的子目录
161+- 延迟数据无固定异常值（如多次运行同一 case 延迟完全相同且不合理）
Aplugins-community/tilelang2ascendc-ops-generator/workflows/task-prompts.md+48-0
@@ -0,0 +1,48 @@
1+# tilelang2ascendc-ops-generator Task Prompts
2+ 
3+调用 Subagent 前必须读取对应 Phase 的完整 prompt 模板。
4+ 
5+---
6+ 
7+## Phase 0-7: 端到端算子开发
8+ 
9+### 调用方式
10+ 
11+调用 `ascend-kernel-developer` Subagent，传入用户需求参数。
12+ 
13+### Prompt 模板
14+ 
15+```
16+你是 ascend-kernel-developer，负责从 PyTorch Model 出发，端到端完成算子设计表达和 AscendC kernel 落地。支持双路径：简单算子走 ops-direct-invoke 工作流（Architect 设计 → Developer 实现 → Reviewer 审查），复杂算子走 TileLang 设计表达 → AscendC 转译。
17+ 
18+## 任务参数
19+ 
20+- NPU 设备: {npu}
21+- 算子描述文件: {op_file}
22+- 输出目录: {output_dir}
23+ 
24+## 执行要求
25+ 
26+按以下 Phase 顺序执行，每完成一个 Phase 汇报状态：
27+ 
28+Phase 0: 参数确认 — 解析参数，设置 ASCEND_RT_VISIBLE_DEVICES={npu}
29+Phase 1: 环境准备 — 创建 {output_dir}/，复制算子文件
30+Phase 2: Case 精简 — 调用 tilelang2ascend-case-simplifier 精简测试用例
31+Phase 3: 设计表达（分支）
32+ ├─ 简单算子: ops-direct-invoke 架构设计 + 设计串讲（DESIGN.md + PLAN.md + WALKTHROUGH.md）
33+ └─ 复杂算子: TileLang 设计 — 调用 tilelang2ascend-tilelang-designer，迭代验证
34+Phase 4: AscendC 生成与验证（分支）
35+ ├─ 简单算子: ops-direct-invoke 开发实现 + 代码审查 + 修复循环（REVIEW.md + 最多3轮修复）
36+ └─ 复杂算子: AscendC 转译 — 调用 tilelang2ascend-translator，迭代验证（最多 3 轮）
37+Phase 5: 性能分析 — 调用 ops-profiling（--compare 模式）
38+Phase 6: 全量验证 — 恢复全量用例，最终验证
39+Phase 7: Trace 记录 — 调用 tilelang2ascend-trace-recorder 生成 trace.md
40+ 
41+## 约束
42+ 
43+- 禁止修改 {output_dir}/ 之外的任何文件
44+- 简单算子 Phase 4 REVIEW.md 修复循环上限 3 轮
45+- 复杂算子 Phase 4 AscendC 验证最多 3 轮迭代
46+- 退化检测必须前置，通过后再执行功能验证
47+- 语言：思考用中文，代码和路径用英文
48+```
Aplugins-community/tilelang2ascendc-ops-generator/workflows/templates/archive_tasks/_kernel_setup_common.py+69-0
@@ -0,0 +1,69 @@
1+#!/usr/bin/env python3
2+# ----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------------------------------------
11+"""Shared kernel setup logic for archive_tasks templates.
12+ 
13+Each task's kernel/setup.py delegates to `kernel_setup(op_name, kernel_dir=__file__)`.
14+"""
15+ 
16+import os
17+import glob
18+import subprocess
19+from pathlib import Path
20+ 
21+import setuptools
22+from setuptools.command.build_ext import build_ext
23+from torch_npu.utils.cpp_extension import NpuExtension
24+ 
25+ 
26+def kernel_setup(op_name, kernel_dir, ext_name=None, description=None):
27+ """Run setuptools.setup() with a standard AscendC kernel build configuration.
28+ 
29+ Args:
30+ op_name: Operator name used as the Python package name.
31+ kernel_dir: Path to the kernel directory (use __file__ from the calling setup.py).
32+ ext_name: NpuExtension module name. Defaults to "{op_name}_ext".
33+ description: Optional one-line description for the package.
34+ """
35+ here = Path(kernel_dir).resolve().parent
36+ build_dir = here / "build"
37+ if ext_name is None:
38+ ext_name = f"{op_name}_ext"
39+ if description is None:
40+ description = f"{op_name} AscendC kernel"
41+ 
42+ class _BuildExt(build_ext):
43+ def run(self):
44+ so_files = glob.glob(str(build_dir / f"{ext_name}*.so"))
45+ if not so_files:
46+ os.makedirs(build_dir, exist_ok=True)
47+ soc_ver = os.environ.get("SOC_VERSION", "Ascend910B2")
48+ ascend = os.environ.get("ASCEND_HOME_PATH", "")
49+ subprocess.check_call(
50+ ["cmake", str(here),
51+ f"-DSOC_VERSION={soc_ver}",
52+ f"-DASCEND_CANN_PACKAGE_PATH={ascend}",
53+ "-DCMAKE_BUILD_TYPE=Release"],
54+ cwd=build_dir)
55+ subprocess.check_call(
56+ ["make", f"-j{os.cpu_count()}"], cwd=build_dir)
57+ 
58+ self.build_lib = str(build_dir)
59+ super().run()
60+ 
61+ setuptools.setup(
62+ name=op_name,
63+ version="0.1.0",
64+ description=description,
65+ ext_modules=[NpuExtension(ext_name, sources=[])],
66+ cmdclass={"build_ext": _BuildExt},
67+ license="BSD 3-Clause",
68+ python_requires=">=3.8",
69+ )
Aplugins-community/tilelang2ascendc-ops-generator/workflows/templates/archive_tasks/_matmul_kernel_common.h+170-0
@@ -0,0 +1,170 @@
1+/*
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef MATMUL_KERNEL_COMMON_H
12+#define MATMUL_KERNEL_COMMON_H
13+ 
14+#include "kernel_operator.h"
15+ 
16+// CRTP base for MatmulKernel templates shared between quant_matmul and matmul_leakyrelu.
17+// Derived must provide accessors: base_m(), base_n(), base_k(), base_mk(), base_kn(), base_mn(),
18+// l1_prefetch().
19+template <typename Derived, typename aType, typename bType, typename cType>
20+class MatmulKernelBase {
21+ static_assert(std::is_same<aType, bType>::value, "aType and bType must be the same type");
22+ static constexpr uint32_t C0 = 32 / sizeof(aType);
23+ 
24+public:
25+ __aicore__ inline MatmulKernelBase() {}
26+ __aicore__ inline void ComputeBlock(const AscendC::GlobalTensor<aType> &aBlock,
27+ const AscendC::GlobalTensor<bType> &bBlock,
28+ const AscendC::GlobalTensor<cType> &cBlock);
29+ 
30+protected:
31+ AscendC::TQue<AscendC::TPosition::A1, 1> inQueueA1;
32+ AscendC::TQue<AscendC::TPosition::A2, 1> inQueueA2;
33+ AscendC::TQue<AscendC::TPosition::B1, 1> inQueueB1;
34+ AscendC::TQue<AscendC::TPosition::B2, 1> inQueueB2;
35+ AscendC::TQue<AscendC::TPosition::CO1, 1> outQueueCO1;
36+ 
37+ uint32_t k_;
38+ uint32_t a_dvalue_, b_dvalue_;
39+ 
40+ __aicore__ inline void CopyA(const AscendC::GlobalTensor<aType> &A, uint32_t kLen);
41+ __aicore__ inline void CopyB(const AscendC::GlobalTensor<bType> &B, uint32_t kLen);
42+ __aicore__ inline void SplitA(const AscendC::LocalTensor<aType> &a1Local,
43+ uint32_t offset, uint32_t colC0Stride);
44+ __aicore__ inline void SplitB(const AscendC::LocalTensor<bType> &b1Local,
45+ uint32_t offset, uint32_t colC0Stride);
46+ __aicore__ inline void Compute(const AscendC::LocalTensor<cType> &c1Local, bool cmatrixInitVal);
47+ __aicore__ inline void CopyOut(const AscendC::GlobalTensor<cType> &C);
48+ 
49+private:
50+ __aicore__ inline Derived &self() { return static_cast<Derived &>(*this); }
51+ __aicore__ inline const Derived &self() const { return static_cast<const Derived &>(*this); }
52+};
53+ 
54+// ---------------------------------------------------------------------------
55+// ComputeBlock
56+// ---------------------------------------------------------------------------
57+template <typename Derived, typename aType, typename bType, typename cType>
58+__aicore__ inline void MatmulKernelBase<Derived, aType, bType, cType>::ComputeBlock(
59+ const AscendC::GlobalTensor<aType> &aBlock,
60+ const AscendC::GlobalTensor<bType> &bBlock,
61+ const AscendC::GlobalTensor<cType> &cBlock)
62+{
63+ AscendC::LocalTensor<cType> c1Local = outQueueCO1.AllocTensor<cType>();
64+ uint32_t kTiles = k_ / self().base_k();
65+ 
66+ for (uint32_t outer = 0; outer < kTiles; outer += self().l1_prefetch()) {
67+ uint32_t count = (kTiles - outer < self().l1_prefetch())
68+ ? (kTiles - outer) : self().l1_prefetch();
69+ uint32_t kLen = count * self().base_k();
70+ 
71+ CopyA(aBlock[outer * self().base_k()], kLen);
72+ CopyB(bBlock[outer * self().base_k() * b_dvalue_], kLen);
73+ 
74+ AscendC::LocalTensor<aType> a1Local = inQueueA1.DeQue<aType>();
75+ AscendC::LocalTensor<bType> b1Local = inQueueB1.DeQue<bType>();
76+ 
77+ for (uint32_t i = 0; i < count; i++) {
78+ SplitA(a1Local, i * self().base_mk(), self().base_m());
79+ SplitB(b1Local, i * self().base_k() * C0, kLen);
80+ Compute(c1Local, (outer + i == 0));
81+ }
82+ 
83+ inQueueA1.FreeTensor(a1Local);
84+ inQueueB1.FreeTensor(b1Local);
85+ }
86+ 
87+ outQueueCO1.EnQue(c1Local);
88+ CopyOut(cBlock);
89+}
90+ 
91+// ---------------------------------------------------------------------------
92+// CopyA
93+// ---------------------------------------------------------------------------
94+template <typename Derived, typename aType, typename bType, typename cType>
95+__aicore__ inline void MatmulKernelBase<Derived, aType, bType, cType>::CopyA(
96+ const AscendC::GlobalTensor<aType> &A, uint32_t kLen)
97+{
98+ auto a1Local = inQueueA1.AllocTensor<aType>();
99+ LoadNdGmToNzL1(a1Local, A, self().base_m(), kLen, a_dvalue_);
100+ inQueueA1.EnQue(a1Local);
101+}
102+ 
103+// ---------------------------------------------------------------------------
104+// CopyB
105+// ---------------------------------------------------------------------------
106+template <typename Derived, typename aType, typename bType, typename cType>
107+__aicore__ inline void MatmulKernelBase<Derived, aType, bType, cType>::CopyB(
108+ const AscendC::GlobalTensor<bType> &B, uint32_t kLen)
109+{
110+ AscendC::LocalTensor<bType> b1Local = inQueueB1.AllocTensor<bType>();
111+ LoadNdGmToNzL1(b1Local, B, kLen, self().base_n(), b_dvalue_);
112+ inQueueB1.EnQue(b1Local);
113+}
114+ 
115+// ---------------------------------------------------------------------------
116+// SplitA
117+// ---------------------------------------------------------------------------
118+template <typename Derived, typename aType, typename bType, typename cType>
119+__aicore__ inline void MatmulKernelBase<Derived, aType, bType, cType>::SplitA(
120+ const AscendC::LocalTensor<aType> &a1Local, uint32_t offset, uint32_t colC0Stride)
121+{
122+ AscendC::LocalTensor<aType> a2Local = inQueueA2.AllocTensor<aType>();
123+ LoadNzL1ToZzL0A(a2Local, a1Local[offset], self().base_m(), self().base_k(), colC0Stride);
124+ inQueueA2.EnQue(a2Local);
125+}
126+ 
127+// ---------------------------------------------------------------------------
128+// SplitB
129+// ---------------------------------------------------------------------------
130+template <typename Derived, typename aType, typename bType, typename cType>
131+__aicore__ inline void MatmulKernelBase<Derived, aType, bType, cType>::SplitB(
132+ const AscendC::LocalTensor<bType> &b1Local, uint32_t offset, uint32_t colC0Stride)
133+{
134+ AscendC::LocalTensor<bType> b2Local = inQueueB2.AllocTensor<bType>();
135+ LoadNzL1ToZnL0B(b2Local, b1Local[offset], self().base_k(), self().base_n(), colC0Stride);
136+ inQueueB2.EnQue(b2Local);
137+}
138+ 
139+// ---------------------------------------------------------------------------
140+// Compute
141+// ---------------------------------------------------------------------------
142+template <typename Derived, typename aType, typename bType, typename cType>
143+__aicore__ inline void MatmulKernelBase<Derived, aType, bType, cType>::Compute(
144+ const AscendC::LocalTensor<cType> &c1Local, bool cmatrixInitVal)
145+{
146+ AscendC::LocalTensor<aType> a2Local = inQueueA2.DeQue<aType>();
147+ AscendC::LocalTensor<bType> b2Local = inQueueB2.DeQue<bType>();
148+ AscendC::MmadParams mmadParams;
149+ mmadParams.m = self().base_m();
150+ mmadParams.n = self().base_n();
151+ mmadParams.k = self().base_k();
152+ mmadParams.cmatrixInitVal = cmatrixInitVal;
153+ AscendC::Mmad(c1Local, a2Local, b2Local, mmadParams);
154+ inQueueA2.FreeTensor(a2Local);
155+ inQueueB2.FreeTensor(b2Local);
156+}
157+ 
158+// ---------------------------------------------------------------------------
159+// CopyOut
160+// ---------------------------------------------------------------------------
161+template <typename Derived, typename aType, typename bType, typename cType>
162+__aicore__ inline void MatmulKernelBase<Derived, aType, bType, cType>::CopyOut(
163+ const AscendC::GlobalTensor<cType> &C)
164+{
165+ auto c1Local = outQueueCO1.DeQue<cType>();
166+ FixpipeNzL0cToNdGm(C, c1Local, self().base_m(), self().base_n());
167+ outQueueCO1.FreeTensor(c1Local);
168+}
169+ 
170+#endif // MATMUL_KERNEL_COMMON_H
Aplugins-community/tilelang2ascendc-ops-generator/workflows/templates/archive_tasks/avg_pool3_d/design/avg_pool3_d_common.py+194-0
@@ -0,0 +1,194 @@
1+#!/usr/bin/env python3
2+# ----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------------------------------------
11+ 
12+"""Shared types and utilities for avg_pool3_d block-level and tile-level designs."""
13+ 
14+from collections import namedtuple
15+from typing import NamedTuple
16+ 
17+import tilelang
18+ 
19+ 
20+pass_configs = {
21+ tilelang.PassConfigKey.TL_ASCEND_AUTO_SYNC: True,
22+ tilelang.PassConfigKey.TL_ASCEND_MEMORY_PLANNING: True,
23+}
24+ 
25+ 
26+class PoolSpatialWindow(NamedTuple):
27+ """Spatial dimension parameters for a pooling kernel window."""
28+ k_h: int
29+ k_w: int
30+ s_h: int
31+ s_w: int
32+ p_h: int
33+ p_w: int
34+ 
35+ 
36+class KernelImpls(NamedTuple):
37+ """Collection of kernel implementation functions for branching dispatch."""
38+ generic_3d: callable
39+ reduce_d: callable
40+ split_c: callable
41+ split_w: callable
42+ multi_w: callable
43+ 
44+ 
45+class PoolParams(NamedTuple):
46+ """Pool kernel parameters for compute_pool_divisor."""
47+ k_d: int
48+ k_h: int
49+ k_w: int
50+ s_d: int
51+ s_h: int
52+ s_w: int
53+ p_d: int
54+ p_h: int
55+ p_w: int
56+ d_dim: int
57+ h_dim: int
58+ w_dim: int
59+ count_include_pad: int
60+ divisor_override: int
61+ pool_size: int
62+ 
63+ 
64+_BlockArgs = namedtuple("_BlockArgs", [
65+ "n_batch", "c_dim", "d_dim", "h_dim", "w_dim", "out_d", "out_h", "out_w",
66+ "k_d", "k_h", "k_w", "s_d", "s_h", "s_w", "p_d", "p_h", "p_w",
67+ "count_include_pad", "divisor_override",
68+ "block_c", "dtype", "split_mode", "split_w_tile_kw", "multi_w_window_w_num",
69+ "in_spatial", "out_spatial", "m_out", "block_m", "m_num",
70+ "used_core_num_m", "tasks_per_core_m", "vec_num", "sub_block_m",
71+ "pool_size", "hw", "split_block_c", "c_num",
72+ "total_blocks_mc", "used_core_num_mc", "tasks_per_core_mc",
73+ "split_w_step", "multi_w_window",
74+])
75+ 
76+ 
77+def parse_block_args(*args):
78+ """Parse variadic avg_pool3_d args and compute block-level tiling parameters."""
79+ a = [*args, 0, "float32", 0, 0, 1]
80+ (
81+ n_batch, c_dim, d_dim, h_dim, w_dim, out_d, out_h, out_w,
82+ k_d, k_h, k_w, s_d, s_h, s_w, p_d, p_h, p_w,
83+ count_include_pad, divisor_override,
84+ block_c, dtype, split_mode, split_w_tile_kw, multi_w_window_w_num,
85+ ) = a[:24]
86+ in_spatial = d_dim * h_dim * w_dim
87+ out_spatial = out_d * out_h * out_w
88+ m_out = n_batch * out_spatial
89+ 
90+ block_m = 0
91+ for candidate in (64, 32, 16, 8, 4, 2):
92+ if candidate <= m_out and m_out % candidate == 0:
93+ block_m = candidate
94+ break
95+ if block_m == 0:
96+ raise ValueError(
97+ f"Unsupported output spatial size: M_out={m_out} is not divisible by any block_M in [2,4,8,16,32,64]"
98+ )
99+ 
100+ m_num = m_out // block_m
101+ if m_num <= 0:
102+ raise ValueError(f"m_num must be positive, got {m_num}")
103+ 
104+ num_physical_cores = 20
105+ used_core_num_m = min(num_physical_cores, m_num)
106+ tasks_per_core_m = (m_num + used_core_num_m - 1) // used_core_num_m
107+ 
108+ vec_num = 2
109+ sub_block_m = block_m // vec_num
110+ 
111+ pool_size = k_d * k_h * k_w
112+ hw = h_dim * w_dim
113+ 
114+ split_block_c = block_c if block_c > 0 else 1
115+ c_num = c_dim // split_block_c
116+ total_blocks_mc = m_num * c_num
117+ used_core_num_mc = min(num_physical_cores, total_blocks_mc)
118+ tasks_per_core_mc = (total_blocks_mc + used_core_num_mc - 1) // used_core_num_mc
119+ 
120+ split_w_step = split_w_tile_kw if split_w_tile_kw > 0 else k_w
121+ multi_w_window = multi_w_window_w_num if multi_w_window_w_num > 0 else 1
122+ 
123+ return _BlockArgs(
124+ n_batch, c_dim, d_dim, h_dim, w_dim, out_d, out_h, out_w,
125+ k_d, k_h, k_w, s_d, s_h, s_w, p_d, p_h, p_w,
126+ count_include_pad, divisor_override,
127+ block_c, dtype, split_mode, split_w_tile_kw, multi_w_window_w_num,
128+ in_spatial, out_spatial, m_out, block_m, m_num,
129+ used_core_num_m, tasks_per_core_m, vec_num, sub_block_m,
130+ pool_size, hw, split_block_c, c_num,
131+ total_blocks_mc, used_core_num_mc, tasks_per_core_mc,
132+ split_w_step, multi_w_window,
133+ )
134+ 
135+ 
136+def decompose_spatial(out_row, out_spatial, out_h, out_w):
137+ """Decompose a flattened output row index into (n_idx, od, oh, ow)."""
138+ n_idx = out_row // out_spatial
139+ out_rem = out_row - n_idx * out_spatial
140+ od = out_rem // (out_h * out_w)
141+ od_rem = out_rem - od * (out_h * out_w)
142+ oh = od_rem // out_w
143+ ow = od_rem - oh * out_w
144+ return n_idx, od, oh, ow
145+ 
146+ 
147+def flatten_5d(*args):
148+ """Flatten a (batch, d, h, w) 5-D coordinate into a row offset."""
149+ n_idx, in_spatial, id_val, hw, h_val, w_dim, w_val = args
150+ return n_idx * in_spatial + id_val * hw + h_val * w_dim + w_val
151+ 
152+ 
153+def compute_pool_divisor(od, oh, ow, params: PoolParams):
154+ """Compute the reciprocal divisor for a pooling window at (od, oh, ow)."""
155+ if params.divisor_override > 0:
156+ return 1.0 / params.divisor_override
157+ if params.count_include_pad > 0:
158+ return 1.0 / params.pool_size
159+ d_begin = max(0, od * params.s_d - params.p_d)
160+ h_begin = max(0, oh * params.s_h - params.p_h)
161+ w_begin = max(0, ow * params.s_w - params.p_w)
162+ d_end = min(params.d_dim, od * params.s_d - params.p_d + params.k_d)
163+ h_end = min(params.h_dim, oh * params.s_h - params.p_h + params.k_h)
164+ w_end = min(params.w_dim, ow * params.s_w - params.p_w + params.k_w)
165+ valid = max(0, d_end - d_begin) * max(0, h_end - h_begin) * max(0, w_end - w_begin)
166+ return 1.0 / valid if valid > 0 else 0.0
167+ 
168+ 
169+def is_reduce_d_fast_path(window: PoolSpatialWindow):
170+ """Return True if only the depth dimension participates in pooling."""
171+ return (window.k_h == 1 and window.k_w == 1
172+ and window.s_h == 1 and window.s_w == 1
173+ and window.p_h == 0 and window.p_w == 0)
174+ 
175+ 
176+def resolve_kernel_mode(window, impls, **kwargs):
177+ """Resolve which kernel implementation to use based on input parameters."""
178+ split_mode = kwargs.get('split_mode', 0)
179+ split_c_ready = kwargs.get('split_c_ready', False)
180+ multi_w_window = kwargs.get('multi_w_window', 1)
181+ default = kwargs.get('default')
182+ if is_reduce_d_fast_path(window):
183+ return impls.reduce_d
184+ if split_mode == 1:
185+ if split_c_ready:
186+ return impls.split_c
187+ raise ValueError("split_c requires divisible channel tiles")
188+ if split_mode == 2:
189+ return impls.split_w
190+ if split_mode == 3:
191+ return impls.multi_w if multi_w_window > 1 else impls.generic_3d
192+ if split_c_ready:
193+ return impls.split_c
194+ return default
Aplugins-community/tilelang2ascendc-ops-generator/workflows/templates/archive_tasks/avg_pool3_d/design/block_level/avg_pool3_d.py+176-0
@@ -0,0 +1,176 @@
1+#!/usr/bin/env python3
2+# ----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------------------------------------
11+ 
12+"""Block-level TileLang design for avg_pool3_d.
13+ 
14+This file mirrors tile-level mode selection and scheduling, but keeps only
15+coarse-grained block partition and Vector-side pipeline skeleton.
16+Fine-grained compute details are intentionally left as TODO comments.
17+"""
18+ 
19+import sys
20+from pathlib import Path
21+ 
22+import tilelang
23+import tilelang.language as T
24+ 
25+# Import shared types and utilities from the common module.
26+_COMMON_DIR = Path(__file__).resolve().parent.parent
27+if str(_COMMON_DIR) not in sys.path:
28+ sys.path.insert(0, str(_COMMON_DIR))
29+from avg_pool3_d_common import (
30+ pass_configs,
31+ PoolSpatialWindow,
32+ KernelImpls,
33+ decompose_spatial,
34+ is_reduce_d_fast_path,
35+ resolve_kernel_mode,
36+ parse_block_args,
37+)
38+ 
39+ 
40+# NOTE: Each @T.prim_func below (generic_3d, reduce_d, split_c, split_w, multi_w)
41+# shares the same pipeline boilerplate (Kernel setup, row/block decomposition,
42+# spatial-dimension flattening). This is intentional: T.lang prim_funcs compile
43+# their body into separate compute graphs and *cannot* share DSL code via a
44+# Python helper. The common structure is listed in full for clarity of each
45+# variant's inner logic.
46+ 
47+ 
48+@tilelang.jit(out_idx=[1], pass_configs=pass_configs)
49+def avg_pool3_d(*args):
50+ (
51+ n_batch, c_dim, d_dim, h_dim, w_dim, out_d, out_h, out_w,
52+ k_d, k_h, k_w, s_d, s_h, s_w, p_d, p_h, p_w,
53+ count_include_pad, divisor_override,
54+ block_c, dtype, split_mode, split_w_tile_kw, multi_w_window_w_num,
55+ in_spatial, out_spatial, m_out, block_m, m_num,
56+ used_core_num_m, tasks_per_core_m, vec_num, sub_block_m,
57+ pool_size, hw, split_block_c, c_num,
58+ total_blocks_mc, used_core_num_mc, tasks_per_core_mc,
59+ split_w_step, multi_w_window,
60+ ) = parse_block_args(*args)
61+ 
62+ def _process_rows(x_in, y_out, row_base, sub_block_m):
63+ for m in T.serial(sub_block_m):
64+ out_row = row_base + m
65+ n_idx, od, oh, ow = decompose_spatial(out_row, out_spatial, out_h, out_w)
66+ _ = (x_in, y_out, n_idx, od, oh, ow)
67+ 
68+ def _process_split_c_rows(x_in, y_out, row_base, c_base, sub_block_m):
69+ for m in T.serial(sub_block_m):
70+ out_row = row_base + m
71+ n_idx, od, oh, ow = decompose_spatial(out_row, out_spatial, out_h, out_w)
72+ _ = (x_in, y_out, c_base, n_idx, od, oh, ow)
73+ 
74+ def _scope_process_rows(x_in, y_out, bx, vid):
75+ with T.Scope("V"):
76+ if bx < m_num:
77+ _process_rows(x_in, y_out, bx * block_m + vid * sub_block_m, sub_block_m)
78+ 
79+ def _scope_process_split_c_rows(*args):
80+ x_in, y_out, task_id, bx, bc, vid = args
81+ with T.Scope("V"):
82+ if task_id < total_blocks_mc:
83+ _process_split_c_rows(x_in, y_out, bx * block_m + vid * sub_block_m,
84+ bc * split_block_c, sub_block_m)
85+ 
86+ def _scope_process_multi_w_rows(x_in, y_out, bx, vid):
87+ with T.Scope("V"):
88+ if bx < m_num:
89+ _process_multi_w_rows(x_in, y_out, bx * block_m + vid * sub_block_m, sub_block_m)
90+ 
91+ def _process_multi_w_aligned(*args):
92+ x_in, y_out, n_idx, od, oh, ow, m, sub_block_m = args
93+ for local_w in T.serial(multi_w_window):
94+ if local_w < sub_block_m - m and local_w < out_w - ow:
95+ _ = (x_in, y_out, n_idx, od, oh, ow + local_w)
96+ 
97+ def _process_multi_w_prefix(*args):
98+ x_in, y_out, n_idx, od, oh, ow, m, sub_block_m, prefix_len = args
99+ for local_w in T.serial(multi_w_window):
100+ if local_w < prefix_len and local_w < sub_block_m - m and local_w < out_w - ow:
101+ _ = (x_in, y_out, n_idx, od, oh, ow + local_w)
102+ 
103+ def _process_multi_w_rows(x_in, y_out, row_base, sub_block_m):
104+ for m in T.serial(sub_block_m):
105+ out_row = row_base + m
106+ n_idx, od, oh, ow = decompose_spatial(out_row, out_spatial, out_h, out_w)
107+ if multi_w_window > 1 and ow % multi_w_window == 0:
108+ _process_multi_w_aligned(x_in, y_out, n_idx, od, oh, ow, m, sub_block_m)
109+ elif m == 0:
110+ prefix_len = multi_w_window - (ow % multi_w_window)
111+ _process_multi_w_prefix(x_in, y_out, n_idx, od, oh, ow, m, sub_block_m, prefix_len)
112+ 
113+ @T.prim_func
114+ def generic_3d(
115+ x_in: T.Tensor((n_batch * in_spatial, c_dim), dtype),
116+ y_out: T.Tensor((m_out, c_dim), dtype),
117+ ):
118+ with T.Kernel(used_core_num_m, is_npu=True) as (cid, vid):
119+ core_idx = cid
120+ for local_idx in T.serial(tasks_per_core_m):
121+ bx = core_idx * tasks_per_core_m + local_idx
122+ _scope_process_rows(x_in, y_out, bx, vid)
123+ 
124+ @T.prim_func
125+ def reduce_d(
126+ x_in: T.Tensor((n_batch * in_spatial, c_dim), dtype),
127+ y_out: T.Tensor((m_out, c_dim), dtype),
128+ ):
129+ with T.Kernel(used_core_num_m, is_npu=True) as (cid, vid):
130+ core_idx = cid
131+ for local_idx in T.serial(tasks_per_core_m):
132+ bx = core_idx * tasks_per_core_m + local_idx
133+ _scope_process_rows(x_in, y_out, bx, vid)
134+ 
135+ @T.prim_func
136+ def split_c(
137+ x_in: T.Tensor((n_batch * in_spatial, c_dim), dtype),
138+ y_out: T.Tensor((m_out, c_dim), dtype),
139+ ):
140+ with T.Kernel(used_core_num_mc, is_npu=True) as (cid, vid):
141+ core_idx = cid
142+ for local_idx in T.serial(tasks_per_core_mc):
143+ task_id = core_idx * tasks_per_core_mc + local_idx
144+ bc = task_id // m_num
145+ bx = task_id % m_num
146+ _scope_process_split_c_rows(x_in, y_out, task_id, bx, bc, vid)
147+ 
148+ @T.prim_func
149+ def split_w(
150+ x_in: T.Tensor((n_batch * in_spatial, c_dim), dtype),
151+ y_out: T.Tensor((m_out, c_dim), dtype),
152+ ):
153+ with T.Kernel(used_core_num_m, is_npu=True) as (cid, vid):
154+ core_idx = cid
155+ for local_idx in T.serial(tasks_per_core_m):
156+ bx = core_idx * tasks_per_core_m + local_idx
157+ _scope_process_rows(x_in, y_out, bx, vid)
158+ 
159+ @T.prim_func
160+ def multi_w(
161+ x_in: T.Tensor((n_batch * in_spatial, c_dim), dtype),
162+ y_out: T.Tensor((m_out, c_dim), dtype),
163+ ):
164+ with T.Kernel(used_core_num_m, is_npu=True) as (cid, vid):
165+ core_idx = cid
166+ for local_idx in T.serial(tasks_per_core_m):
167+ bx = core_idx * tasks_per_core_m + local_idx
168+ _scope_process_multi_w_rows(x_in, y_out, bx, vid)
169+ 
170+ split_c_ready = block_c > 0 and c_dim % block_c == 0
171+ window = PoolSpatialWindow(k_h=k_h, k_w=k_w, s_h=s_h, s_w=s_w, p_h=p_h, p_w=p_w)
172+ impls = KernelImpls(generic_3d=generic_3d, reduce_d=reduce_d,
173+ split_c=split_c, split_w=split_w, multi_w=multi_w)
174+ return resolve_kernel_mode(window, impls,
175+ split_mode=split_mode, split_c_ready=split_c_ready,
176+ multi_w_window=multi_w_window)
Aplugins-community/tilelang2ascendc-ops-generator/workflows/templates/archive_tasks/avg_pool3_d/design/tile_level/avg_pool3_d.py+269-0
@@ -0,0 +1,269 @@
1+#!/usr/bin/env python3
2+# ----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------------------------------------
11+ 
12+import sys
13+from pathlib import Path
14+ 
15+import tilelang
16+import tilelang.language as T
17+ 
18+# Tile-level path setup differs from block-level.
19+_p = Path(__file__).resolve().parent.parent
20+if str(_p) not in sys.path:
21+ sys.path.insert(0, str(_p))
22+import avg_pool3_d_common as _common
23+pass_configs = _common.pass_configs
24+PoolSpatialWindow = _common.PoolSpatialWindow
25+KernelImpls = _common.KernelImpls
26+decompose_spatial = _common.decompose_spatial
27+flatten_5d = _common.flatten_5d
28+compute_pool_divisor = _common.compute_pool_divisor
29+is_reduce_d_fast_path = _common.is_reduce_d_fast_path
30+resolve_kernel_mode = _common.resolve_kernel_mode
31+parse_block_args = _common.parse_block_args
32+ 
33+# NOTE: Each @T.prim_func below (generic_3d, reduce_d, split_c, split_w, multi_w)
34+# shares the same pipeline boilerplate (Kernel setup, UB allocation, row/block
35+# decomposition, spatial-dimension flattening). This is intentional: T.lang
36+# prim_funcs compile their body into separate compute graphs and *cannot* share
37+# DSL code via a Python helper. The common structure is listed in full for
38+# clarity of each variant's inner logic.
39+ 
40+ 
41+@tilelang.jit(out_idx=[1], pass_configs=pass_configs)
42+def avg_pool3_d(*args):
43+ a = parse_block_args(*args)
44+ (
45+ n_batch, c_dim, d_dim, h_dim, w_dim, out_d, out_h, out_w,
46+ k_d, k_h, k_w, s_d, s_h, s_w, p_d, p_h, p_w,
47+ count_include_pad, divisor_override,
48+ block_c, dtype, split_mode, split_w_tile_kw, multi_w_window_w_num,
49+ in_spatial, out_spatial, m_out, block_m, m_num,
50+ _, _, vec_num, sub_block_m,
51+ pool_size, hw, split_block_c, c_num,
52+ _, _, _,
53+ split_w_step, multi_w_window,
54+ ) = a
55+ 
56+ def _process_scope(*args, mode="generic"):
57+ if mode == "split_c":
58+ x_in, y_out, acc_ub, inp_ub, out_ub, row_base, c_base = args
59+ else:
60+ x_in, y_out, acc_ub, inp_ub, out_ub, row_base = args
61+ c_base = 0
62+ with T.Scope("V"):
63+ for m in T.serial(sub_block_m):
64+ _process_scope_row(
65+ x_in, y_out, acc_ub, inp_ub, out_ub,
66+ row_base + m, c_base, mode,
67+ )
68+ 
69+ def _process_scope_row(*args):
70+ x_in, y_out, acc_ub, inp_ub, out_ub, out_row, c_base, mode = args
71+ n_idx, od, oh, ow = decompose_spatial(out_row, out_spatial, out_h, out_w)
72+ T.tile.fill(acc_ub, T.float32(0.0))
73+ for kd in T.serial(k_d):
74+ if mode == "reduce_d":
75+ _accumulate_reduce_d_slice(x_in, acc_ub, inp_ub, n_idx, in_spatial, hw, od, oh, ow, kd)
76+ elif mode == "split_w":
77+ _accumulate_split_w_slice(x_in, acc_ub, inp_ub, n_idx, in_spatial, hw, od, oh, ow, kd)
78+ else:
79+ _accumulate_depth_slice(x_in, acc_ub, inp_ub, n_idx, in_spatial, hw, od, oh, ow, kd, col=c_base)
80+ if divisor_override > 0:
81+ T.tile.mul(out_ub, acc_ub, T.float32(1.0 / divisor_override))
82+ else:
83+ T.tile.mul(out_ub, acc_ub, T.float32(1.0 / pool_size))
84+ T.copy(out_ub, y_out[out_row, c_base])
85+ 
86+ @T.prim_func
87+ def generic_3d(
88+ x_in: T.Tensor((n_batch * in_spatial, c_dim), dtype),
89+ y_out: T.Tensor((m_out, c_dim), dtype),
90+ ):
91+ with T.Kernel(m_num, is_npu=True) as (cid, vid):
92+ bx = cid
93+ row_base = bx * block_m + vid * sub_block_m
94+ acc_ub = T.alloc_ub((1, c_dim), dtype)
95+ inp_ub = T.alloc_ub((1, c_dim), dtype)
96+ out_ub = T.alloc_ub((1, c_dim), dtype)
97+ _process_scope(x_in, y_out, acc_ub, inp_ub, out_ub, row_base, mode="generic")
98+ 
99+ @T.prim_func
100+ def reduce_d(
101+ x_in: T.Tensor((n_batch * in_spatial, c_dim), dtype),
102+ y_out: T.Tensor((m_out, c_dim), dtype),
103+ ):
104+ with T.Kernel(m_num, is_npu=True) as (cid, vid):
105+ bx = cid
106+ row_base = bx * block_m + vid * sub_block_m
107+ acc_ub = T.alloc_ub((1, c_dim), dtype)
108+ inp_ub = T.alloc_ub((1, c_dim), dtype)
109+ out_ub = T.alloc_ub((1, c_dim), dtype)
110+ _process_scope(x_in, y_out, acc_ub, inp_ub, out_ub, row_base, mode="reduce_d")
111+ 
112+ @T.prim_func
113+ def split_c(
114+ x_in: T.Tensor((n_batch * in_spatial, c_dim), dtype),
115+ y_out: T.Tensor((m_out, c_dim), dtype),
116+ ):
117+ with T.Kernel(m_num * c_num, is_npu=True) as (cid, vid):
118+ bx = cid % m_num
119+ bc = cid // m_num
120+ row_base = bx * block_m + vid * sub_block_m
121+ c_base = bc * split_block_c
122+ acc_ub = T.alloc_ub((1, split_block_c), dtype)
123+ inp_ub = T.alloc_ub((1, split_block_c), dtype)
124+ out_ub = T.alloc_ub((1, split_block_c), dtype)
125+ _process_scope(x_in, y_out, acc_ub, inp_ub, out_ub, row_base, c_base, mode="split_c")
126+ 
127+ @T.prim_func
128+ def split_w(
129+ x_in: T.Tensor((n_batch * in_spatial, c_dim), dtype),
130+ y_out: T.Tensor((m_out, c_dim), dtype),
131+ ):
132+ with T.Kernel(m_num, is_npu=True) as (cid, vid):
133+ bx = cid
134+ row_base = bx * block_m + vid * sub_block_m
135+ acc_ub = T.alloc_ub((1, c_dim), dtype)
136+ inp_ub = T.alloc_ub((1, c_dim), dtype)
137+ out_ub = T.alloc_ub((1, c_dim), dtype)
138+ _process_scope(x_in, y_out, acc_ub, inp_ub, out_ub, row_base, mode="split_w")
139+ 
140+ def _accumulate_reduce_d_slice(*args, col=0):
141+ x_in, acc_ub, inp_ub, n_idx, in_spatial, hw, od, oh, ow, kd = args
142+ id_val = od * s_d - p_d + kd
143+ if 0 <= id_val < d_dim:
144+ in_row = flatten_5d(n_idx, in_spatial, id_val, hw, oh, w_dim, ow)
145+ T.copy(x_in[in_row, col], inp_ub)
146+ T.tile.add(acc_ub, acc_ub, inp_ub)
147+ 
148+ def _accumulate_width_slice(*args, col=0):
149+ x_in, acc_ub, inp_ub, n_idx, in_spatial, hw, id_val, ih_val, cur_ow = args
150+ for kw in T.serial(k_w):
151+ iw_val = cur_ow * s_w - p_w + kw
152+ if 0 <= iw_val < w_dim:
153+ in_row = flatten_5d(
154+ n_idx, in_spatial,
155+ id_val, hw, ih_val,
156+ w_dim, iw_val)
157+ T.copy(x_in[in_row, col], inp_ub)
158+ T.tile.add(acc_ub, acc_ub, inp_ub)
159+ 
160+ def _accumulate_depth_slice(*args, col=0):
161+ x_in, acc_ub, inp_ub, n_idx, in_spatial, hw, od, oh, cur_ow, kd = args
162+ id_val = od * s_d - p_d + kd
163+ if 0 <= id_val < d_dim:
164+ for kh in T.serial(k_h):
165+ ih_val = oh * s_h - p_h + kh
166+ if 0 <= ih_val < h_dim:
167+ _accumulate_width_slice(x_in, acc_ub, inp_ub, n_idx, in_spatial, hw,
168+ id_val, ih_val, cur_ow, col=col)
169+ 
170+ def _accumulate_split_w_element(*args):
171+ x_in, acc_ub, inp_ub, n_idx, in_spatial, hw, id_val, ih_val, iw_val = args
172+ if 0 <= iw_val < w_dim:
173+ in_row = flatten_5d(
174+ n_idx, in_spatial,
175+ id_val, hw, ih_val,
176+ w_dim, iw_val)
177+ T.copy(x_in[in_row, 0], inp_ub)
178+ T.tile.add(acc_ub, acc_ub, inp_ub)
179+ 
180+ def _accumulate_split_w_chunk(*args):
181+ x_in, acc_ub, inp_ub, n_idx, in_spatial, hw, id_val, ih_val, ow = args
182+ for kw_base in T.serial((k_w + split_w_step - 1) // split_w_step):
183+ for kw_local in T.serial(split_w_step):
184+ kw = kw_base * split_w_step + kw_local
185+ if kw < k_w:
186+ iw_val = ow * s_w - p_w + kw
187+ _accumulate_split_w_element(x_in, acc_ub, inp_ub, n_idx, in_spatial, hw,
188+ id_val, ih_val, iw_val)
189+ 
190+ def _accumulate_split_w_slice(*args):
191+ x_in, acc_ub, inp_ub, n_idx, in_spatial, hw, od, oh, ow, kd = args
192+ id_val = od * s_d - p_d + kd
193+ if 0 <= id_val < d_dim:
194+ for kh in T.serial(k_h):
195+ ih_val = oh * s_h - p_h + kh
196+ if 0 <= ih_val < h_dim:
197+ _accumulate_split_w_chunk(x_in, acc_ub, inp_ub, n_idx, in_spatial, hw,
198+ id_val, ih_val, ow)
199+ 
200+ def _accumulate_window(*args):
201+ x_in, acc_ub, inp_ub, out_ub, y_out, out_row, local_w, n_idx, in_spatial, hw, od, oh, cur_ow = args
202+ T.tile.fill(acc_ub, T.float32(0.0))
203+ for kd in T.serial(k_d):
204+ _accumulate_depth_slice(x_in, acc_ub, inp_ub, n_idx, in_spatial, hw, od, oh, cur_ow, kd)
205+ divisor = compute_pool_divisor(
206+ od, oh, cur_ow, k_d, k_h, k_w, s_d, s_h, s_w, p_d, p_h, p_w,
207+ d_dim, h_dim, w_dim, count_include_pad, divisor_override, pool_size)
208+ if divisor > 0:
209+ T.tile.mul(out_ub, acc_ub, T.float32(divisor))
210+ else:
211+ T.tile.fill(out_ub, T.float32(0.0))
212+ T.copy(out_ub, y_out[out_row + local_w, 0])
213+ 
214+ def _process_multi_w_aligned(*args):
215+ x_in, y_out, acc_ub, inp_ub, out_ub, out_row, n_idx, in_spatial, hw, od, oh, ow, m, sub_block_m = args
216+ for local_w in T.serial(multi_w_window):
217+ if local_w < sub_block_m - m and local_w < out_w - ow:
218+ _accumulate_window(
219+ x_in, acc_ub, inp_ub, out_ub, y_out, out_row, local_w,
220+ n_idx, in_spatial, hw, od, oh, ow + local_w)
221+ 
222+ def _process_multi_w_prefix(*args):
223+ x_in, y_out, acc_ub, inp_ub, out_ub, out_row, n_idx, in_spatial, hw, od, oh, ow, m, sub_block_m, prefix_len = (
224+ args
225+ )
226+ for local_w in T.serial(multi_w_window):
227+ if local_w < prefix_len and local_w < sub_block_m - m and local_w < out_w - ow:
228+ _accumulate_window(
229+ x_in, acc_ub, inp_ub, out_ub, y_out, out_row, local_w,
230+ n_idx, in_spatial, hw, od, oh, ow + local_w)
231+ 
232+ def _process_multi_w_branch(*args):
233+ x_in, y_out, acc_ub, inp_ub, out_ub, out_row, n_idx, in_spatial, hw, od, oh, ow, m, sub_block_m = args
234+ if multi_w_window > 1 and ow % multi_w_window == 0:
235+ _process_multi_w_aligned(x_in, y_out, acc_ub, inp_ub, out_ub, out_row,
236+ n_idx, in_spatial, hw, od, oh, ow, m, sub_block_m)
237+ elif m == 0:
238+ prefix_len = multi_w_window - (ow % multi_w_window)
239+ _process_multi_w_prefix(x_in, y_out, acc_ub, inp_ub, out_ub, out_row,
240+ n_idx, in_spatial, hw, od, oh, ow, m, sub_block_m, prefix_len)
241+ 
242+ def _process_multi_w_row(*args):
243+ x_in, y_out, acc_ub, inp_ub, out_ub, row_base = args
244+ with T.Scope("V"):
245+ for m in T.serial(sub_block_m):
246+ out_row = row_base + m
247+ n_idx, od, oh, ow = decompose_spatial(out_row, out_spatial, out_h, out_w)
248+ _process_multi_w_branch(x_in, y_out, acc_ub, inp_ub, out_ub, out_row,
249+ n_idx, in_spatial, hw, od, oh, ow, m, sub_block_m)
250+ 
251+ @T.prim_func
252+ def multi_w(
253+ x_in: T.Tensor((n_batch * in_spatial, c_dim), dtype),
254+ y_out: T.Tensor((m_out, c_dim), dtype),
255+ ):
256+ with T.Kernel(m_num, is_npu=True) as (cid, vid):
257+ bx = cid
258+ row_base = bx * block_m + vid * sub_block_m
259+ acc_ub = T.alloc_ub((1, c_dim), dtype)
260+ inp_ub = T.alloc_ub((1, c_dim), dtype)
261+ out_ub = T.alloc_ub((1, c_dim), dtype)
262+ _process_multi_w_row(x_in, y_out, acc_ub, inp_ub, out_ub, row_base)
263+ 
264+ split_c_ready = block_c > 0 and c_dim % block_c == 0
265+ window = PoolSpatialWindow(k_h=k_h, k_w=k_w, s_h=s_h, s_w=s_w, p_h=p_h, p_w=p_w)
266+ impls = KernelImpls(generic_3d=generic_3d, reduce_d=reduce_d,
267+ split_c=split_c, split_w=split_w, multi_w=multi_w)
268+ return resolve_kernel_mode(split_mode, split_c_ready, multi_w_window,
269+ window, impls, default=impls.generic_3d)
Aplugins-community/tilelang2ascendc-ops-generator/workflows/templates/archive_tasks/avg_pool3_d/kernel/CMakeLists.txt+94-0
@@ -0,0 +1,94 @@
1+# set the library output dir to the python dir for wheel package build
2+set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${PROJECT_SOURCE_DIR}/python/ascend_kernel/ascend_kernel/lib)
3+ 
4+# host side files
5+FILE(GLOB OP_SRCS
6+ ${PROJECT_OP_SRC_BASE}/register.cpp
7+ ${PROJECT_OP_SRC_BASE}/ops/helloworld/op_host/helloworld.cpp
8+ ${PROJECT_OP_SRC_BASE}/ops/abs/op_host/abs.cpp
9+ ${PROJECT_OP_SRC_BASE}/ops/cumsum/op_host/cumsum.cpp
10+ ${PROJECT_OP_SRC_BASE}/ops/histc/op_host/histc.cpp
11+ ${PROJECT_OP_SRC_BASE}/ops/sum/op_host/sum.cpp
12+ ${PROJECT_OP_SRC_BASE}/ops/avg_pool3_d/op_host/avg_pool3_d.cpp
13+ )
14+if(BUILD_CATLASS_MODULE)
15+ list(APPEND OP_SRCS
16+ ${PROJECT_OP_SRC_BASE}/ops/catlass/op_host/catlass_matmul_basic.cpp
17+ )
18+endif()
19+ 
20+# set the so name
21+set(OP_PLUGIN_NAME ascend_kernel)
22+ 
23+# kernel side files without workspace
24+ascendc_library(no_workspace_kernel STATIC
25+ ${PROJECT_OP_SRC_BASE}/ops/helloworld/op_kernel/kernel_helloworld.cpp
26+ ${PROJECT_OP_SRC_BASE}/ops/abs/op_kernel/abs.cpp
27+ ${PROJECT_OP_SRC_BASE}/ops/cumsum/op_kernel/cumsum.cpp
28+ ${PROJECT_OP_SRC_BASE}/ops/histc/op_kernel/histc.cpp
29+ ${PROJECT_OP_SRC_BASE}/ops/sum/op_kernel/sum.cpp
30+ ${PROJECT_OP_SRC_BASE}/ops/avg_pool3_d/op_kernel/avg_pool3_d_generic.cpp
31+ ${PROJECT_OP_SRC_BASE}/ops/avg_pool3_d/op_kernel/avg_pool3_d_reduce_d.cpp
32+ ${PROJECT_OP_SRC_BASE}/ops/avg_pool3_d/op_kernel/avg_pool3_d_split_c.cpp
33+ ${PROJECT_OP_SRC_BASE}/ops/avg_pool3_d/op_kernel/avg_pool3_d_split_w.cpp
34+ ${PROJECT_OP_SRC_BASE}/ops/avg_pool3_d/op_kernel/avg_pool3_d_multi_w.cpp
35+ )
36+ 
37+# # kernel side files with workspace
38+# set(WORKSPACE_KERNEL_SRCS
39+# ${PROJECT_OP_SRC_BASE}/ops/mla_preprocess/op_kernel/mla_preprocess_kernel.cpp
40+# )
41+# if(BUILD_CATLASS_MODULE)
42+# list(APPEND WORKSPACE_KERNEL_SRCS
43+# ${PROJECT_OP_SRC_BASE}/ops/catlass/op_kernel/catlass_matmul_basic_kernel.cpp
44+# )
45+# endif()
46+# ascendc_library(workspace_kernel STATIC ${WORKSPACE_KERNEL_SRCS})
47+# if(BUILD_CATLASS_MODULE)
48+# ascendc_include_directories(workspace_kernel PRIVATE
49+# ${CATLASS_DIR}/include
50+# )
51+# endif()
52+ 
53+# ascendc_compile_definitions(workspace_kernel PRIVATE
54+# -DHAVE_WORKSPACE
55+# -DHAVE_TILING
56+# )
57+ 
58+# create shared library libascend_kernel.so
59+add_library(${OP_PLUGIN_NAME} SHARED ${OP_SRCS})
60+ 
61+target_link_libraries(${OP_PLUGIN_NAME} PRIVATE
62+ # workspace_kernel
63+ no_workspace_kernel
64+ torch_npu
65+ ascendcl
66+ tiling_api
67+ nnopbase
68+ opapi
69+ register
70+ platform
71+ ascendalog
72+ dl
73+)
74+ 
75+target_link_directories(${OP_PLUGIN_NAME} PRIVATE
76+ ${TORCH_DIR}/lib
77+ ${TORCH_NPU_DIR}/lib
78+)
79+ 
80+target_include_directories(${OP_PLUGIN_NAME} PRIVATE
81+ ${CMAKE_CURRENT_SOURCE_DIR}/utils
82+ ${PROJECT_OP_SRC_BASE}/utils
83+ ${PROJECT_SOURCE_DIR}/include
84+ ${PROJECT_OP_SRC_BASE}/ops/avg_pool3_d/op_kernel
85+ ${TORCH_DIR}/include
86+ ${TORCH_DIR}/include/torch/csrc/api/include
87+ ${TORCH_NPU_DIR}/include/third_party/acl/inc
88+ ${TORCH_NPU_DIR}/include/third_party/hccl/inc
89+ ${TORCH_NPU_DIR}/include
90+ ${PYTHON_INCLUDE_DIR}
91+ ${ASCEND_INCLUDE_DIR}/external
92+ ${ASCEND_INCLUDE_DIR}/experiment/platform
93+ ${ASCEND_INCLUDE_DIR}/experiment/runtime
94+)
Aplugins-community/tilelang2ascendc-ops-generator/workflows/templates/archive_tasks/avg_pool3_d/kernel/build.sh+134-0
@@ -0,0 +1,134 @@
1+#!/bin/bash
2+# ----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------------------------------------
11+ 
12+ 
13+set -e
14+ 
15+BUILD_KERNELS_MODULE="ON"
16+DEBUG_MODE="OFF"
17+ 
18+while getopts ":a:hd" opt; do
19+ case ${opt} in
20+ a )
21+ BUILD_KERNELS_MODULE="OFF"
22+ case "$OPTARG" in
23+ kernels )
24+ BUILD_KERNELS_MODULE="ON"
25+ ;;
26+ * )
27+ echo "Error: Invalid Value"
28+ echo "Allowed value: kernels"
29+ exit 1
30+ ;;
31+ esac
32+ ;;
33+ d )
34+ DEBUG_MODE="ON"
35+ ;;
36+ h )
37+ echo "Use './build.sh' build all modules."
38+ echo "Use './build.sh -a <target>' to build specific parts of the project."
39+ echo " <target> can be:"
40+ echo " kernels Only build ascend_kernel."
41+ exit 1
42+ ;;
43+ \? )
44+ echo "Error: unknown flag: -$OPTARG" 1>&2
45+ echo "Run './build.sh -h' for more information."
46+ exit 1
47+ ;;
48+ : )
49+ echo "Error: -$OPTARG requires a value" 1>&2
50+ echo "Run './build.sh -h' for more information."
51+ exit 1
52+ ;;
53+ esac
54+done
55+ 
56+shift $((OPTIND -1))
57+ 
58+ 
59+export DEBUG_MODE=$DEBUG_MODE
60+ 
61+SOC_VERSION="${1:-Ascend910_9382}"
62+ 
63+ 
64+### Get Current CANN Toolkit Installation Path
65+_CANN_TOOLKIT_INSTALL_PATH=$(cat /etc/Ascend/ascend_cann_install.info | grep "Toolkit_InstallPath" | awk -F'=' '{print $2}')
66+source ${_CANN_TOOLKIT_INSTALL_PATH}/set_env.sh
67+echo -e "\e[1;32mDetected CANN Toolkit Installation Path: ${_CANN_TOOLKIT_INSTALL_PATH}\e[0m"
68+echo -e "\e[1;33mDouble Checking Environment Variables:\e[0m"
69+echo -e "\e[1;32mASCEND_HOME_PATH: ${ASCEND_HOME_PATH}\e[0m"
70+echo -e "\e[1;32mASCEND_TOOLKIT_HOME: ${ASCEND_TOOLKIT_HOME}\e[0m"
71+ 
72+ 
73+ASCEND_INCLUDE_DIR=${ASCEND_TOOLKIT_HOME}/$(arch)-linux/include
74+CURRENT_DIR=$(pwd)
75+PROJECT_ROOT=$(dirname "$CURRENT_DIR")
76+VERSION="1.0.0"
77+OUTPUT_DIR=$CURRENT_DIR/output
78+mkdir -p $OUTPUT_DIR
79+echo "outpath: ${OUTPUT_DIR}"
80+ 
81+COMPILE_OPTIONS=""
82+ 
83+function build_kernels()
84+{
85+ CMAKE_DIR=""
86+ BUILD_DIR="build"
87+ 
88+ cd "$CMAKE_DIR" || exit
89+ 
90+ rm -rf $BUILD_DIR
91+ mkdir -p $BUILD_DIR
92+ 
93+ cmake $COMPILE_OPTIONS \
94+ -DCMAKE_INSTALL_PREFIX="$OUTPUT_DIR" \
95+ -DASCEND_HOME_PATH=$ASCEND_HOME_PATH \
96+ -DASCEND_INCLUDE_DIR=$ASCEND_INCLUDE_DIR \
97+ -DSOC_VERSION=$SOC_VERSION \
98+ -DBUILD_KERNELS_MODULE=$BUILD_KERNELS_MODULE \
99+ -B "$BUILD_DIR" \
100+ -S .
101+ 
102+ cmake --build "$BUILD_DIR" --target install -j 16
103+ cd -
104+}
105+ 
106+function make_ascend_kernel_package()
107+{
108+ cd python/ascend_kernel || exit
109+ 
110+ rm -rf "$CURRENT_DIR"/python/ascend_kernel/dist
111+ cp -v "${CURRENT_DIR}/config.ini" "${CURRENT_DIR}/python/ascend_kernel/ascend_kernel/"
112+ python3 setup.py clean --all
113+ python3 setup.py bdist_wheel
114+ mv -v "$CURRENT_DIR"/python/ascend_kernel/dist/ascend_kernel*.whl ${OUTPUT_DIR}/
115+ rm -rf "$CURRENT_DIR"/python/ascend_kernel/dist
116+ cd -
117+}
118+ 
119+function main()
120+{
121+ 
122+ build_kernels
123+ if pip3 show wheel;then
124+ echo "wheel has been installed"
125+ else
126+ pip3 install wheel==0.45.1
127+ fi
128+ if [[ "$BUILD_KERNELS_MODULE" == "ON" ]]; then
129+ make_ascend_kernel_package
130+ fi
131+ 
132+}
133+ 
134+main
Aplugins-community/tilelang2ascendc-ops-generator/workflows/templates/archive_tasks/avg_pool3_d/kernel/design.md+54-0
@@ -0,0 +1,54 @@
1+# avg_pool3_d 算子设计文档
2+ 
3+## 算子概述
4+ 
5+`avg_pool3_d` 是 3D 平均池化算子，支持 4D/5D 输入张量 (NCDHW/NDHWC)。
6+ 
7+## 实现变体
8+ 
9+本算子提供 5 种 AscendC kernel 实现，根据输入特征自动选择：
10+ 
11+| 实现 | 触发条件 | 说明 |
12+|------|---------|------|
13+| `generic` | 默认 | 通用实现，逐点计算池化窗口平均值 |
14+| `reduce_d` | kH=1, kW=1, sH=1, sW=1 | 仅对 D 维度做 reduce，优化一维情况 |
15+| `split_c` | tilelang_mode=split_c | 沿 C 维度切分，利用向量化并行 |
16+| `split_w` | tilelang_mode=split_w | 沿 W 维度切分，优化宽方向大 kernel |
17+| `multi_w` | tilelang_mode=multi_w | 多窗口并行，提升 W 方向吞吐 |
18+ 
19+## 目录结构
20+ 
21+```
22+csrc/ops/avg_pool3_d/
23+├── CMakeLists.txt
24+├── design.md
25+├── op_host/
26+│ └── avg_pool3_d.cpp # Host 侧封装：参数处理、tiling、kernel 分发
27+└── op_kernel/
28+ ├── avg_pool3_d_generic.cpp
29+ ├── avg_pool3_d_generic_kernel.h
30+ ├── avg_pool3_d_reduce_d.cpp
31+ ├── avg_pool3_d_reduce_d_kernel.h
32+ ├── avg_pool3_d_split_c.cpp
33+ ├── avg_pool3_d_split_c_kernel.h
34+ ├── avg_pool3_d_split_w.cpp
35+ ├── avg_pool3_d_split_w_kernel.h
36+ ├── avg_pool3_d_multi_w.cpp
37+ ├── avg_pool3_d_multi_w_kernel.h
38+ ├── avg_pool3_d_tiling.h
39+ └── kernel_common.h
40+```
41+ 
42+## Host 侧逻辑
43+ 
44+1. **输出尺寸计算**：参考 PyTorch `avg_pool3d` 公式，支持 `ceil_mode`
45+2. **数据格式转换**：NCDHW -> NDHWC flat (n*d*h*w, c)
46+3. **Tiling 构建**：填充 `AvgPool3DKernelTiling` 结构，包含所有几何参数和调度参数
47+4. **Kernel 分发**：根据 `kH/kW/sH/sW/pH/pW` 选择最优实现变体
48+5. **结果回写**：NDHWC -> NCDHW，4D 输入去除 channel 维度
49+ 
50+## 注册信息
51+ 
52+- `ops.h`：`at::Tensor avg_pool3d(...)`
53+- `register.cpp`：`torch.ops.npu.avg_pool3d`
54+- `csrc/CMakeLists.txt`：已添加 host/kernel 源文件，移除旧 ACLNN 封装