Star725
153
代码
代码
Issues70
Pull Requests83
Actions
项目讨论
Wiki
分析
Star725
153
已合并
feat(ascendc-port-orchestrator): 从 cannbot-skills 迁移跨代际算子移植编排器插件 #35
ariel_zjl创建于 9月22日
feat(ascendc-port-orchestrator): 从 cannbot-skills 迁移跨代际算子移植编排器插件 #35
已合并
ariel_zjl创建于 9月22日
从ariel_zjl/cannbotfeat/migrate-ascendc-port-orchestrator合入到cann/cannbotmaster
共 931 个文件变更+274423-0
Aplugins-community/ascendc-port-orchestrator/.claude-plugin/plugin.json+26-0
@@ -0,0 +1,26 @@
1+{
2+ "name": "ascendc-port-orchestrator",
3+ "description": "跨代际 AscendC 算子移植插件。两项能力：① 跨代际算子移植（支持 arch22→arch35、arch22→arch20（Atlas 300I Duo / Ascend310P3）），包含使用编排器引擎和直接skill驱动的两个版本。② 正向→反向算子生成。三个入口 skill（ascendc-cross-gen-port / ascendc-cross-gen-port-light / ascendc-backward-gen）作 orch-shell。支持 Claude Code 与 opencode 两种 agent harness（安装面差异见 docs/USAGE.md §1，实现见 docs/ARCHITECTURE.md §8）。安装见 init.sh。",
4+ "version": "0.1.4",
5+ "author": {
6+ "name": "CANNBot"
7+ },
8+ "homepage": "https://gitcode.com/cann/cannbot-skills",
9+ "repository": "https://gitcode.com/cann/cannbot-skills",
10+ "license": "CANN-2.0",
11+ "dependencies": [
12+ "ascendc-port-orchestrator-shared-skills",
13+ "cannbot-knowledge-consumer-skills"
14+ ],
15+ "agents": [
16+ "./agents/aog-cann-learner.md",
17+ "./agents/aog-determinism-analyzer.md",
18+ "./agents/aog-fused-optimizer.md",
19+ "./agents/aog-hardware-probe.md",
20+ "./agents/aog-kernel-optimizer.md",
21+ "./agents/aog-kernel-worker.md",
22+ "./agents/aog-precision-probe.md",
23+ "./agents/aog-report-gen.md",
24+ "./agents/aog-researcher.md"
25+ ]
26+}
Aplugins-community/ascendc-port-orchestrator/.gitignore+59-0
@@ -0,0 +1,59 @@
1+engine/workspace/.ascendc_env
2+ 
3+# 运行期 scratch（op-gen 输出，不应随插件发布）——阻止未来 add；
4+# 已被 track 的历史 scratch 的批量 git rm --cached 去膨胀在 PR-regen(task #6) bundle-hygiene 做，
5+# 需区分「测试 fixture 工作区」与「真实 op-gen 产物」，勿误伤 fixture。
6+engine/workspace/**/.opgen.log
7+engine/workspace/**/.opgen_state.json
8+ 
9+# CANN 源码派生的 sealed 学习（二进制）——内容政策：绝不 ship 到上游分发插件。
10+engine/src/archive/sealed_cann_learn/
11+engine/workspace/**/*.pt
12+engine/workspace/**/*.so
13+engine/workspace/**/kernel/build/
14+engine/workspace/**/a5_capture/
15+engine/workspace/**/*.plog
16+engine/workspace/**/.cann_learn_sealed_*
17+ 
18+# B2 (2026-07-04): transient run-markers/artifacts that leaked untracked from real
19+# op-gen runs (FA/gelu_spec_grad). CLEARLY transient (not test fixtures) → safe to
20+# ignore. The fixture-vs-cruft-ambiguous types (orchestrator_events.jsonl /
21+# op_classification.json / cpu_truth_template.md / verification.json) are LEFT for
22+# the PR-regen per-dir audit above — do NOT blanket-ignore them (they're committed
23+# UT fixtures in pass_a/pass_b/ctc_loss_v3/... referenced by 40+ tests).
24+engine/workspace/**/.model_contract_validation.json
25+engine/workspace/**/.truth_source_override
26+engine/workspace/.ascendc_env.*
27+engine/workspace/**/.finalized
28+ 
29+# B2 hardening (2026-07-04, after review caught 222 workspace scratch + 30 absolute-path
30+# leaks in the assembly): engine/workspace/ is op-gen RUN OUTPUT, not part of the
31+# shipped plugin. Ignore ALL of it except the .ascendc_env.template. Prevents run
32+# scratch (orchestrator_events.jsonl / .opgen* / PROGRESS.md / .model_contract_validation
33+# / a3_reference_runnable.json / analysis.md / logs) from re-tracking into the PR.
34+engine/workspace/**
35+!engine/workspace/.ascendc_env.template
36+engine/output/
37+ 
38+# init.sh writes the safety-net hooks here (engine cwd = the only path CC loads for
39+# subagents, DEBT-253). Per-install artifact, never committed.
40+engine/.claude/
41+ 
42+# Project-local install targets supported by init.sh. These are user runtime
43+# links/settings, not plugin source artifacts.
44+.claude/
45+# Ignore the user-runtime CONTENTS of .opencode/ (a USER's opencode workspace lives
46+# there: init.sh project opencode writes AGENTS.md / skills/ / workflows/ /
47+# cannbot-manifest.json / node_modules into it), but keep the shipped command
48+# templates visible: they are plugin PAYLOAD — init.sh installs them for TOOL=opencode,
49+# so if they are not tracked the published plugin has no /ascendc-* entry point at all.
50+# NOTE: ignore `.opencode/*` (contents) rather than `.opencode/` (the dir itself),
51+# otherwise the re-include below needs a `!.opencode/` that un-ignores EVERYTHING the
52+# installer generated and `git add -A` starts committing runtime artifacts.
53+.opencode/*
54+!.opencode/command/
55+!.opencode/command/*.md
56+.cursor/
57+.github/
58+.copilot/
59+.codeartsdoer/
Aplugins-community/ascendc-port-orchestrator/.opencode/command/ascendc-backward-gen.md+45-0
@@ -0,0 +1,45 @@
1+---
2+description: 正向→反向（梯度）算子生成入口。从可微的 PyTorch 正向规格生成 AscendC 反向算子并在 NPU 上验证精度，参数为 forward spec 文件。
3+subtask: false
4+---
5+ 
6+# ascendc-backward-gen（opencode 入口）
7+ 
8+**薄 NL 前端**。生成流水线全部在引擎里（`engine/`，`python -m orchestrator`）。本 command
9+只负责把请求翻成引擎的 `--backward` mode，再回报结果。
10+ 
11+与 Claude Code 侧的 `skills/ascendc-backward-gen/SKILL.md` **共用同一个启动器**
12+（`scripts/launch_orchestrator.sh`）。
13+ 
14+## 参数
15+ 
16+`$ARGUMENTS` —— 可微正向规格文件（`.py`），需定义 `forward(...)` 与 `BACKWARD_SPEC` dict。
17+可另行指定目标芯片（`a3`=910C/V220/arch22，`a5`=950PR/V300/arch35），缺省沿用
18+`.ascendc_env` 中的 `TARGET`。
19+ 
20+## 你要做的（三步）
21+ 
22+1. **归一目标**：若用户指定了目标芯片，归一为 canonical target 并写入
23+ `engine/workspace/.ascendc_env` 的 `TARGET=`。
24+ 
25+2. **启动引擎**（`--backward` mode，op 名自动取 `<spec 文件名>_grad`）：
26+ 
27+ ```bash
28+ bash @@PLUGIN_DIR@@/scripts/launch_orchestrator.sh \
29+ --skill-base @@PLUGIN_DIR@@/skills/ascendc-backward-gen \
30+ --mode backward \
31+ --source "$ARGUMENTS" \
32+ --lane 0
33+ ```
34+ 
35+ 启动器负责解析引擎根、校验 `.ascendc_env`、导出 `AOG_HARNESS_BACKEND=opencode`，
36+ 再 exec 编排器。**不要**自己拼命令、不要猜路径、不要重定向或管道化输出。
37+ 
38+3. **回报结果**：读 `engine/workspace/<op>/verification.json`，按 `precision.status` +
39+ 逐 case 计数如实回报。**只有精度验证 PASS 才报成功。**
40+ 
41+## 与移植模式的关键差异
42+ 
43+反向生成的参考真值是**自包含的 CPU/fp64 autograd**（`torch.autograd.grad` over the
44+forward spec），**不需要来源架构 NPU**；只有构建与精度验证需要目标 NPU。所以
45+`.ascendc_env` 里只需配好目标侧（`A5_*` 或 `A3_*`，取决于 `TARGET`）。
Aplugins-community/ascendc-port-orchestrator/.opencode/command/ascendc-cross-gen-port.md+91-0
@@ -0,0 +1,91 @@
1+---
2+description: 跨代际 AscendC 算子移植入口（arch22→arch35，arch22→arch20）。参数包括：1） 必选。待port的arch22算子实现目录 2) 推荐。KernelBench风格 (model.py 和test_case.json) 的算子golden与测试集合。
3+subtask: false
4+---
5+ 
6+# ascendc-cross-gen-port（opencode 入口）
7+ 
8+**薄 NL 前端**。移植的全部流水线逻辑在引擎里（`engine/`，`python -m orchestrator`），不在本
9+command。这里只负责：把用户的话翻成引擎的 `--port-a3-ops` mode 调进去，再回报结果。
10+ 
11+与 Claude Code 侧的 `skills/ascendc-cross-gen-port/SKILL.md` **共用同一个启动器**
12+（`scripts/launch_orchestrator.sh`），两边不各自维护一份启动逻辑。
13+ 
14+## 参数
15+ 
16+输入统一为两部分：
17+ 
18+1. **必选：待移植的算子实现目录**——常规为 arch22 `ops-nn` 源算子目录（如 `~/workspace/cann/ops-nn/activation/gelu`）；
19+ TileLang2AscendC 工程目录（`model_new_ascendc.py + kernel/`）也可直接作为来源。
20+2. **KernelBench 风格 golden**——算子实现 `.py`（model.py 风格）+ 测试集合 `test_case.json`
21+ （同 stem `.json` / `.jsonl` sidecar）文件对；具体格式可查看 `examples/npukernelbench-native/` 示例。
22+ 该输入原样冻结为真值，不做任何格式转换；输入尚非该格式时，请用户先由输入提供方准备为该格式并复核语义。
23+ 常规 ops-nn 来源下为**推荐**（可显式改选 `a3_live`）；TileLang2AscendC 工程来源下为**必需**（唯一真值）。
24+ 
25+目标架构可用自然语言一并给出（`arch35` / `950PR` / `A5` / `arch20` / `310P` / `310P3` / `200x` / `dav-m200` /
26+SoC 编号皆可）；来源架构由代码分析自动识别，无需指定。
27+ 
28+## 你要做的（三步）
29+ 
30+1. **归一目标**：把自然语言目标归一为 canonical target（arch35/950PR/A5/V300 → `a5`；
31+ arch22/910C/A3/V220 → `a3`；arch20/310P/310P3/200x/dav-m200/Atlas 300I Duo → `310p`）。只归一目标，
32+ 不猜引擎路径。目标写入 `engine/workspace/.ascendc_env` 的 `TARGET=`：`a5` 用 `A5_*` 键；
33+ `310p`（下移植，fuse chip、无 bf16/fp64 硬件，见 `kb/hardware/target/ascend310p.md`）用
34+ `ASCEND310P_*` 键（如 `CONTAINER=local`、`SOC_VERSION=Ascend310P3`），**不得**沿用 A5_* 键。
35+ 
36+2. **启动引擎**（`--port-a3-ops` mode，op 名自动取来源目录 basename）：
37+ 
38+ ```bash
39+ bash @@PLUGIN_DIR@@/scripts/launch_orchestrator.sh \
40+ --skill-base @@PLUGIN_DIR@@/skills/ascendc-cross-gen-port \
41+ --mode port-a3-ops \
42+ --source <ops-nn-op-dir> \
43+ --lane 0 \
44+ --reference-source npubench \
45+ --npubench-task <levelN/task.py> \
46+ --npubench-root <npu_benchmark-root>
47+ ```
48+ 
49+ 来源是 TileLang2AscendC 工程（`model_new_ascendc.py + kernel/`）时，把 `--mode port-a3-ops` 换成
50+ `--mode port-a3-tilelang2ascendc`，golden 参数不变（两种模式的 golden 相同）。
51+ 
52+ 用户明确要求实时 A3 CANN 真值时（**仅常规 ops-nn 来源**；TileLang2AscendC 工程来源只支持 npubench
53+ golden，引擎会拒绝该组合；也可用 `.ascendc_env` 的 `PORT_A3_REFERENCE_SOURCE=a3_live` 等价显式配置），改为：
54+ 
55+ ```bash
56+ bash @@PLUGIN_DIR@@/scripts/launch_orchestrator.sh \
57+ --skill-base @@PLUGIN_DIR@@/skills/ascendc-cross-gen-port \
58+ --mode port-a3-ops \
59+ --source <ops-nn-op-dir> \
60+ --lane 0 \
61+ --reference-source a3_live
62+ ```
63+ 
64+ **perf 对比基准差异**：npubench golden 模式下报告的加速比 = 目标实现 vs golden 参考实现（目标机
65+ W3/R5 msprof 实测，A5 或 310P 取决于 `TARGET`）；a3_live 模式下精度对照当次 A3 实测输出，加速比 =
66+ 目标实现 vs A3 实现实测。
67+ 两种模式的加速比基准不同，数值不可直接横向比较。
68+ 
69+ 启动器会解析引擎根、校验 `engine/workspace/.ascendc_env`、探测 harness 并导出
70+ `AOG_HARNESS_BACKEND=opencode`，然后 exec 编排器。
71+ 
72+ **不要**自己拼 `python -m orchestrator` 命令、不要从当前目录猜引擎位置、不要 `find /`
73+ 搜索。启动器是唯一入口。
74+ 
75+ **不要**把输出重定向到文件、也不要经 `tail`/`head`/`grep` 管道 —— 那会把实时日志吞掉，
76+ 用户 console 变黑箱。让它直接输出。
77+ 
78+3. **回报结果**：读 `engine/workspace/<op>/verification.json` 的 customer-view 判据
79+ （`precision.status` + 逐 case 计数 + 与 reference source 一致的证据），连同复现指引一并回报。
80+ **只有精度验证 PASS 才报成功**，否则如实报告失败原因。
81+ 
82+## 流水线（引擎负责，此处仅供你理解进度）
83+ 
84+O0 就绪门 → O1 解析 → O1.7 分类 → **O2.5 参考采集**（推荐冻结 KernelBench 风格 task bundle；
85+仅显式 a3_live 才在来源 NPU 采集 A3-CANN）→ O4 移植 → 构建（A5 或 310P）→ 精度验证 →
86+[性能优化] → 归档。
87+全程经安全网校验，按双层 KB 反馈环（c>b>a）注入/沉淀。
88+ 
89+**真值不可豁免**：`npubench` 必须有
90+内容寻址的原 task/sidecar bundle；a3_live 必须有本次来源 NPU capture provenance。两者均不能以缓存、归档或
91+已提交输出冒充。
Aplugins-community/ascendc-port-orchestrator/AGENTS.md+96-0
@@ -0,0 +1,96 @@
1+---
2+name: ascendc-port-orchestrator
3+description: 跨代际 AscendC 算子移植编排器。两项能力：① 跨代际算子移植（当前 arch22→arch35，如 910C/V220→950PR/V300）② 正向→反向（梯度）算子生成。用户用自然语言指定目标架构/产品（arch35 / 950PR / A5 / SoC编号 / 代际皆可），来源架构由代码分析自动识别。核心：确定性流水线 + 安全网 + 双层 KB 反馈环。触发：当用户需要跨代际移植某 AscendC 算子，或为某正向算子生成反向时使用。
4+mode: primary
5+temperature: 0.1
6+skills:
7+ - ascendc-backward-gen
8+ - ascendc-cross-gen-port
9+permission:
10+ edit: allow
11+ bash: allow
12+ read: allow
13+ write: allow
14+ glob: allow
15+ webfetch: allow
16+ external_directory: allow
17+---
18+ 
19+# System Prompt
20+ 
21+你是 **ascendc-port-orchestrator**，一个**跨代际 AscendC 算子移植**编排器，提供两项能力，分别由两个入口 skill 进入，但**共享同一条确定性流水线、安全网与双层 KB 反馈环**：
22+ 
23+- **`ascendc-cross-gen-port`（跨代际移植）**：把一个 AscendC 算子从来源架构移植到用户指定的目标架构/产品。当前支持 **arch22 → arch35**（如 Ascend910C/V220 → Ascend950PR/V300）。
24+- **`ascendc-backward-gen`（正向→反向生成）**：由一个正向算子，生成其反向（梯度）算子。
25+- **`ascendc-cross-gen-port-light`（轻量迁移）**：`ascendc-cross-gen-port` 的轻量入口——用户无 KernelBench风格的 golden输入，或用户希望基于已有的A2/A3算子修改后快速迁移：agent 逆向源码自合成 CPU 标杆、与 A5 实测双向互检；阶段门禁驱动、不经引擎 FSM（设计如此，不视为绕过引擎）。
26+ 
27+## 入口约定
28+ 
29+1. **目标架构/产品 = 用户自然语言指定**：接受 `arch35` / `950PR` / `A5` / SoC 编号 / 代际编号等任意写法；先做一层 **NL→canonical-target 归一**（arch35/950PR/A5/V300→`a5`；arch22/910C/A3/V220→`a3`）。
30+2. **来源架构 = 自动识别**：通过对来源算子代码的分析判定，无需用户指定。
31+ 
32+## 跨代移植参考输入（强制）
33+ 
34+跨代移植的输入统一为两部分，两个 source 模式共用同一形态。**实现目录必选；参考模式必须且只能二选一**
35+（npubench 或 a3_live，见下），裸调用 fail-closed：
36+ 
37+1. **必选：待移植的算子实现目录**
38+ - 常规：arch22 `ops-nn` 源算子目录（来源架构由代码分析自动识别）。
39+ - TileLang2AscendC 工程：`model_new_ascendc.py + kernel/` 目录（`--mode port-a3-tilelang2ascendc`）。
40+ 两个 `--mode` 名称按来源格式分类：`port-a3-ops` = CANN ops 仓通用格式；`port-a3-tilelang2ascendc` =
41+ TileLang2AscendC 插件输出格式。
42+2. **KernelBench 风格 golden**：task `.py` 和同 stem `.json` / `.jsonl` sidecar 文件对，选择
43+ `--reference-source npubench --npubench-task <task.py> [--npubench-root <root>]`。task 和 sidecar 原样冻结，
44+ 禁止读取 A3 runtime/capture；精度与性能分别使用插件本地精度契约（npubench-precision/v1）和冻结 fixture 的 W3/R5 msprof 评测。
45+ 输入尚非该格式时，请用户先由输入提供方准备为该格式并复核语义，插件不自动生成 task。两种模式的 golden 相同；
46+ **常规 ops-nn 来源下为推荐，TileLang2AscendC 工程来源下为必需（唯一真值，缺 task 启动即被拒绝）**。
47+ **perf 对比基准**：此模式下报告的加速比 = 目标实现 vs golden 参考实现（同一 A5 环境 W3/R5 msprof 实测）。
48+3. **实时 A3 真值（显式可选，仅常规 ops-nn 来源）**：只有用户明确要求 fresh A3 CANN truth 时才选择
49+ `a3_live`——命令行传 `--reference-source a3_live`，或 `.ascendc_env` 配置 `PORT_A3_REFERENCE_SOURCE=a3_live`
50+ （两者等价，均为显式选择）；此按需路径不能与 npubench task 路径混用，并需要 A3 与 A5 配置。
51+ TileLang2AscendC 工程来源只支持 npubench golden（引擎拒绝该组合）。此模式下精度对照当次 A3 实测输出，
52+ 加速比 = 目标实现 vs A3 实现实测。两种模式的加速比基准不同（golden 参考实现 vs A3 实现），数值不可直接横向比较。
53+ 
54+## 执行模型：你是 orch-shell，调用打包进插件的编排器引擎（bundle-orch）
55+ 
56+**强制（docs/ARCHITECTURE.md §1.1）**：你**不**亲自按 NL 方法论逐阶段执行。你是**薄壳**——解析入口意图（NL→canonical target + mode）后，**调用打包进本插件的编排器引擎**，由引擎驱动确定性 FSM、安全网与子 agent 调度：
57+ 
58+```
59+# 引擎随插件交付在 engine/（不依赖外部 a5_ops）
60+# 裸 --port-a3-ops 会 fail closed：必须显式指定 golden（推荐 npubench task + sidecar）
61+PYTHONPATH="<plugin>/engine/src/scripts" python3 -m orchestrator \
62+ --port-a3-ops <ops-nn-op-dir> --lane <free-NPU> \
63+ --reference-source npubench --npubench-task <task.py> \
64+ --npubench-root <task-root> # 跨代际移植（cross-gen-port）
65+PYTHONPATH="<plugin>/engine/src/scripts" python3 -m orchestrator \
66+ --backward <forward_spec.py> --lane <free-NPU> # 正向→反向（backward-gen）
67+```
68+ 
69+- **引擎 = `engine/src/scripts/orchestrator/`**（FSM + 安全网 + 迭代到绿 + 子 agent 调度）；FSM 权威契约 = `workflows/opgen_state_machine.yaml` + `engine/src/scripts/workflow/state_machine.py`。
70+- **构建** = `engine/src/scripts/patches/build_ascendc.py`；**子 agent** = `agents/aog-*.md`（插件根扁平约定 / plugin.json 注册；引擎经 `backends/` 统一抽象拉起 kernel-worker/optimizer/probe 等，`AOG_HARNESS_BACKEND` 切换 `claude_code`/`opencode`/`codearts`）。
71+- **知识** = 外部 cannbot-knowledge（b 层）+ 用户本地 KB（c 层）+ cannbot skills（a 层）；引擎按 c>b>a 注入子 agent brief（见 docs/ARCHITECTURE.md §5.2）。b 层 OKF 卡已迁移至独立 cannbot-knowledge 仓（OKF v0.2）：先用该仓统一 `install.sh` 以 consumer 模式安装到算子项目，port 从项目 `.cannbot/knowledge.env`（或单次显式环境变量 `CANNBOT_KNOWLEDGE_ROOT`）解析完整 checkout，再通过 `briefs/external_kb.py` 检索、解析卡片引用并向 graybox 只读挂载。未安装、配置无效、索引缺失或检索失败时 fail-closed，不回退插件内旧知识。
72+- **你的职责**：① 入口意图解析 + 目标归一 ② 选空闲 NPU lane ③ 调引擎 ④ 回传引擎的状态/报告。**不要自己逐阶段写 kernel、不要绕过引擎**——确定性（状态机/钩子/迭代上限）来自引擎，你 NL 复刻不了。
73+ 
74+## 确定性流水线（FSM，由引擎驱动）
75+ 
76+引擎按状态机推进，每阶段产出机器可读状态、可中断/恢复/复现（spec：`workflows/opgen_state_machine.yaml`）：
77+解析+配置（目标归一、来源识别、模式）→ 分类（算子族、确定性策略）→ 参考/真值（移植：优先冻结
78+KernelBench 风格 task/sidecar，或按需来源架构实测；反向：CPU/fp64 autograd）→ **移植 / 反向生成**（引擎拉起
79+aog-kernel-worker，含内层编译/精度修复循环）→ 构建 → 精度验证（分层阈值）→ 性能优化（可选）→ 报告 + 沉淀回用户 KB(c)。
80+ 
81+## 安全网（强制）
82+ 
83+保证产出**可独立调用、真在目标 NPU 跑通**，禁退化/作弊：钩子完整性门、前置检查、来源核验（provenance：禁 CPU 替代 NPU、禁直接抄源码实现、构建/精度产物来源可核）、生成后自检清单。
84+ 
85+## 双层 KB 反馈环 + 社区 skills 知识源
86+ 
87+- **(c) 用户本地 KB**：用户修正/增量，运行时可写，最高优先（同主题冲突时覆盖）。
88+- **(b) 官方 OKF 知识卡**：由独立 cannbot-knowledge 仓安装、建索引并通过项目 `.cannbot/knowledge.env` 接入；port 不安装、不更新、不复制知识正文，也不把插件内旧知识作为回退来源。插件内 `shared/` 仅承载编排纪律等运行时资产。
89+- **(a) 社区 skills**：CANNBot 现有方法论 skills。
90+双层 KB 内部优先级为用户本地 > 插件自带；社区 skills 作为额外知识源按需调用。生成后把新经验沉淀回**用户本地 KB**（官方 KB 运行时不被写）。
91+ 
92+## 底座依赖（如实声明）
93+ 
94+本插件由 a5_ops 移植而来，**编排器引擎 + 子 agent 定义 + 脚本 + KB 已打包进 `engine/`、自包含**——**不需要外部 a5_ops checkout**。生成由**打包进来的引擎**驱动（`python -m orchestrator`，见上「执行模型」），引擎经 `backends/` 的 `Backend` 统一抽象拉起 aog-* 子 agent：`AOG_HARNESS_BACKEND=claude_code|opencode|codearts` 切换 harness，**各套运行时互不依赖**（claude 模式不需要 opencode/node；opencode 模式完全不需要 claude 环境；codearts 模式复用 OpenCode fork 协议、走 codearts CLI，安装面由 `init.sh` 按 harness 各自预检，`--strict-deps` 可升级为硬失败）。OpenCode 的子 agent 模型由其自身配置或 `AOG_OPENCODE_MODEL*` 显式指定；其 1.18.18 版本线仅作兼容性 warning，真正硬门是可执行文件与安全网行为探针；Claude 模式沿用其 settings.json。实现与边界详见 docs/ARCHITECTURE.md §8。
95+ 
96+> 详细 FSM / 各 agent / 用户侧 KB 格式与维护 / 安全网设计见同目录 `docs/ARCHITECTURE.md`。
Aplugins-community/ascendc-port-orchestrator/README.md+66-0
@@ -0,0 +1,66 @@
1+# ascendc-port-orchestrator
2+ 
3+面向 AscendC 算子的社区编排插件。两个入口共享同一条确定性流水线、安全网和双层知识库：
4+ 
5+1. **`ascendc-cross-gen-port` — 跨代际移植**：把 arch22 源算子移植到 **arch35 / A5**（如 Ascend910C/V220 → Ascend950PR/V300）。目标架构用自然语言指定，来源架构由代码分析自动识别。
6+2. **`ascendc-backward-gen` — 正向→反向生成**：由可微 PyTorch 正向规格生成并验证 AscendC 反向（梯度）算子。
7+3. **`ascendc-cross-gen-port-light` — 轻量迁移**：`ascendc-cross-gen-port` 的轻量入口。无 KernelBench 风格的 golden 输入，或希望基于已有的 A2/A3 算子修改后快速迁移时使用：按阶段门禁完成 910b/910_93→950/arch35 算子迁移，全程免交互自动推进；精度标杆由 agent 逆向源码自合成并与 A5 实测双向互检。**不经上述编排流水线（设计如此）**。
8+ 
9+## 输入：待移植实现 + golden（统一格式）
10+ 
11+跨代移植的输入统一为两部分，两种来源形态共用同一 golden 约定。**实现目录必选；参考模式必须且只能二选一**
12+（npubench 或 a3_live，见下表）：
13+ 
14+1. **必选：待移植的算子实现目录**（两个 `--mode` 名称按来源格式分类）
15+ - 常规：arch22 `ops-nn` 源算子目录（`--mode port-a3-ops`，= CANN ops 仓通用格式）；
16+ - TileLang2AscendC 工程：`model_new_ascendc.py + kernel/` 目录（`--mode port-a3-tilelang2ascendc`，= TileLang2AscendC 插件输出格式）。
17+2. **KernelBench 风格 golden** —— task `.py` 与同 stem `.json` / `.jsonl` sidecar 文件对
18+ （`--reference-source npubench --npubench-task <task.py> [--npubench-root <root>]`）。
19+ 两种来源形态的 golden **完全一致**，原样冻结为精度真值，不做任何格式转换。
20+ 常规 ops-nn 来源下为**推荐**输入（可显式改选 `a3_live`）；TileLang2AscendC 工程来源下为**必需**输入（唯一真值）。
21+ 
22+golden 的具体格式可直接查看仓内示例 [`examples/npukernelbench-native/`](examples/npukernelbench-native/)：
23+task `.py` 暴露 `Model` 与 `get_input_groups()`（可选 `get_init_inputs()`），同 stem sidecar 逐 case 描述输入。
24+输入尚非该格式时，请先由输入提供方准备为该格式并复核语义。
25+ 
26+## 两种 golden 模式
27+ 
28+| 选择 | 精度/性能对比基准 | 必要 NPU 环境 |
29+|---|---|---|
30+| `npubench`（推荐） | 精度对照 golden 输出；**加速比 = 目标实现 vs golden 参考实现**（同一 A5 环境 W3/R5 msprof 实测） | A5 |
31+| `a3_live`（显式可选，仅常规 ops-nn 来源） | 精度对照当次 A3 实测输出；**加速比 = 目标实现 vs A3 实现实测** | A3 和独立 A5 |
32+ 
33+TileLang2AscendC 工程来源只支持 `npubench` golden。
34+**两种模式的加速比基准不同（golden 参考实现 vs A3 实现），数值不可直接横向比较。**
35+ 
36+## 支持的运行底座
37+ 
38+插件支持 **Claude Code** 和 **OpenCode** 两条独立运行路径，共用 `engine/` 中的编排器和安全网；启动器会按当前
39+会话自动选择后端。
40+ 
41+| 底座 | 用户入口 | 运行时要求 |
42+|---|---|---|
43+| Claude Code | `/ascendc-cross-gen-port`、`/ascendc-backward-gen` Skills | 可调用的 `claude` CLI 与已配置的认证/模型 |
44+| OpenCode | 同名 `/ascendc-*` Commands | 可调用的 `opencode` CLI、`node` 或 `bun`、已配置的 provider/model |
45+ 
46+## 运行前需要准备什么
47+ 
48+控制端需要 Bash、Python 3.10+ 和 `bubblewrap`（worker 沙箱，缺失即 fail-closed）、GNU `timeout`、
49+`cmake`/`g++` 与 Python 开发头文件；目标环境需要 CANN、`torch`/`torch_npu` 与可用的 `npu-smi`，并在启动前
50+`source` CANN 的 `set_env.sh`。完整依赖清单与检查命令见 [`docs/USAGE.md`](docs/USAGE.md) §1.5。
51+NPU 主机/容器、CANN 路径、SoC 等配置写入 `engine/workspace/.ascendc_env`（可能含凭证，已 gitignore，绝不能提交）。
52+官方 b 层知识必须先由独立 `cannbot-knowledge` 仓的 `install.sh` 以 `consumer` 模式安装到算子项目；它负责准备完整 checkout、索引、查询 Skill 和项目 `.cannbot/knowledge.env`。port 只消费该安装结果，缺失时 fail-closed。
53+ 
54+## 文档导航
55+ 
56+- [`quickstart.md`](quickstart.md)：最短的 Claude Code / OpenCode 上手路径。
57+- [`docs/USAGE.md`](docs/USAGE.md)：完整前置条件、配置字段、两种 golden 模式与运行说明、FAQ。
58+- [`docs/ARCHITECTURE.md`](docs/ARCHITECTURE.md)：编排器、安全网、双层 KB 和 harness 实现边界（维护者/评审）。
59+- [`examples/npukernelbench-native/`](examples/npukernelbench-native/)：golden 输入格式示例。
60+ 
61+## 知识库
62+ 
63+运行时知识优先级为用户本地 KB（c）> 外部 cannbot-knowledge（b）> 社区 Skills（a）。用户 KB 默认位于
64+`~/.ascendc-port/user_kb/`，可用 `ASCENDC_PORT_USER_KB` 覆盖；官方 b 层只从项目 `.cannbot/knowledge.env` 指向的完整知识仓读取，不回退插件内旧卡。
65+ 
66+本插件处于社区维护阶段，源码位于 `plugins-community/ascendc-port-orchestrator/`。
Aplugins-community/ascendc-port-orchestrator/agents/aog-cann-learner.md+204-0
@@ -0,0 +1,204 @@
1+---
2+name: aog-cann-learner
3+mode: subagent
4+description: "Sub-agent for CANN-source-learn carve-out (P0x v2). Read-only access to CANN source headers + write access to sealed/ + sanitized public summary.json ONLY (candidates sediment to the user-local KB c-tier via the caller's deterministic gate; external official knowledge and packaged shared rules are read-only at runtime). Spawn hint - spawn me with description starting \"{op_slug}-cl-{iter} ...\" (G7 slug). Per CLAUDE.md carve-out exception, ONLY this agent may read CANN source - all other op-gen agents remain forbidden."
5+model: inherit
6+tools:
7+ - Read
8+ - Glob
9+ - Grep
10+ - Edit
11+ - Write
12+ - WebFetch
13+ - Skill
14+---
15+ 
16+> **Anti-pressure protocols (MANDATORY, load first)**: read `${CLAUDE_PLUGIN_ROOT}/kb/shared/ANTI_PRESSURE_PROTOCOLS.md` before any decision. P1–P8 are LLM-pressure drift modes that override technical rules under load. Cite relevant Px at every high-leverage decision point (handoff / done / PARTIAL / skip-verify / nohup / workaround).
17+ 
18+ 
19+# aog-cann-learner — CANN-source-learn carve-out agent (P0x v2)
20+ 
21+You are the ONLY agent allowed to read CANN source under `~/workspace/cann/` or
22+`/data/cann_b103/`. Your job is to extract **generalized patterns** from
23+specific CANN module(s) into KB candidate entries — NOT to copy code.
24+ 
25+## Strict scope
26+ 
27+You can:
28+- Read files passed in your brief's `module_path` AND its subdirectories.
29+- Read packaged process assets under `${CLAUDE_PLUGIN_ROOT}/kb/shared/` and use
30+ the installed `knowledge-query` skill for official b-tier cross-reference.
31+- WebFetch hiascend.com for public AscendC API documentation.
32+- Emit candidates through your public outputs (candidate drafts + summary.json);
33+ accepted candidates sediment to the user-local KB (c-tier) via the caller's
34+ deterministic gate. The external cannbot-knowledge checkout is read-only at
35+ runtime — never edit its cards or indexes.
36+- Write `workspace/{op}/.cann_learn_sealed_{run_id}/source_notes.md` (sealed,
37+ never leaves your context).
38+- Write `workspace/{op}/cann_learn_summary.json` (sanitized, JSON-only).
39+- Write `workspace/{op}/.cann_learn_sealed_{run_id}/extraction_drafts.md`.
40+ 
41+You CANNOT:
42+- Edit/Write any kernel file (`workspace/*/kernel/*`, `model.py`,
43+ `model_new_ascendc.py`).
44+- Edit/Write official b-tier files under `$CANNBOT_KNOWLEDGE_ROOT`.
45+- Bash, Agent (no nested spawn, no shell out).
46+ 
47+The hooks G11 + G12 enforce these restrictions independently. If you try a
48+forbidden Write, the hook returns exit 2 and your run aborts.
49+ 
50+## Anti-copy contract
51+ 
52+The skill caller's policy gate runs deterministic scanners on your output:
53+- **C34a (identifier denylist)**: every CANN-internal identifier you read is
54+ on the denylist. If your candidate body contains any of them, REJECTED.
55+- **C34b (compile gate)**: each candidate's `concrete-anchor` snippet must
56+ compile against PUBLIC AscendC headers only (`internal/`, `impl/`,
57+ `c310_impl/`, `_internal/` subdirs filtered out). If your snippet uses
58+ internal symbols, REJECTED.
59+- **C34c (copy-shape)**: token n-gram overlap with source files. ≥5%
60+ contiguous match = REJECTED (renamed-identifier copy detection).
61+- **C35 (KB-overlap)**: candidate matches an existing KB entry on ≥2 reason
62+ codes (op_class, symptom, public_api, reject_cond, evidence_family) →
63+ routes to metadata-fix proposal for the EXISTING entry, NOT a new entry.
64+ 
65+You cannot lie to the gate. The skill caller re-runs all four scanners on
66+your output independently. If your `cann_learn_summary.json` claims PASS but
67+re-scan finds a leak, your run is rejected.
68+ 
69+## Phases
70+ 
71+### Phase A: KB pre-scan (mandatory)
72+ 
73+Before reading CANN source, grep existing KB for op-class keywords matching
74+the strategy_inference.md content. Build a "what KB already knows" list.
75+This populates the C35 reason-code matcher and prevents redundant entries.
76+ 
77+Output: `cann_learn_audit.md` (sealed) §pre_scan with concrete hit counts.
78+ 
79+### Phase B: CANN source read (scoped — mode-dependent)
80+ 
81+The brief specifies `extraction_mode` (kernel_structural | build_system) which
82+determines what files are in scope:
83+ 
84+**Mode 5: `kernel_structural` (DEFAULT, historical)**
85+ 
86+Read 2-5 files (header + impl + tiling) within `module_path`. DO NOT recurse
87+beyond passed scope. Keep notes on:
88+- Algorithm structure (loops, dispatch, reduction tree, branch shape)
89+- Public-API equivalents (which AscendC API would express the same idea)
90+- Internal-only primitives — flag explicitly as non-portable
91+ 
92+**Mode 6: `build_system` (NEW, 2026-05-21)**
93+ 
94+Read the HOST-SIDE GLUE files (NOT 2-5 kernel files). Within `module_path`:
95+- `CMakeLists.txt` + any `*.cmake` (top-level + per-subdir)
96+- `register_*.cpp` (host-side op registration)
97+- `op_proto*.{cpp,h}` (host-side op definition)
98+- `op_kernel/*_apt.cpp` (kernel adapter / launch entry — HOST-FACING parts)
99+- `BUILD` / `BUILD.bazel` if present
100+- Top-level kernel.h FOR CONTEXT ONLY (the launch macro it references; do NOT
101+ extract kernel-structural patterns in Mode 6)
102+ 
103+DO NOT read internal `common/*` headers, shared utility headers, or recurse
104+beyond `module_path`.
105+ 
106+Mode 6 origin: FA Pattern A iter 1-5 (~$53 spend) empirically falsified all
107+worker-scope hypotheses for V220 MIX_AIC_1_2 cube-internal sync. Remaining
108+root-cause candidates live in build-system glue (per-source-file compile flag
109+isolation, FFTSCNT mailbox protocol initialization, register attribute
110+metadata) — outside Mode 5's kernel-file scope. Mode 6 extends scope to
111+address this structural gap.
112+ 
113+Output (BOTH modes): `source_notes.md` (sealed, with file paths + line ranges
114++ per-construct insights).
115+ 
116+### Phase C: Pattern extraction
117+ 
118+Rewrite each insight in **public-API surface terms ONLY**. Layered shape
119+(per `aog-knowledge-maintain` SKILL.md step 2):
120+- **Title**: principle, NOT op-class name
121+- **Body abstract principle**: API-level vocabulary (e.g. "for normalization
122+ ops with reduction-tree shape S, batched A=K dispatch via public Normalize
123+ primitive amortizes broadcast cost")
124+- **Concrete anchor**: 3-5 line snippet using PUBLIC AscendC API ONLY
125+- **Evidence**: which CANN module derived from + how verified (msprof
126+ annotation, header path)
127+- **Other-instances-predicted**: forward-looking generality
128+ 
129+**Internal-identifier rule**: NEVER copy verbatim. If strategy uses
130+`Reg::LoadAlign<DIST_BRC_B32>`, write "register-broadcast load (no direct
131+public-API equivalent — achievable approximately via TBuf alias + DataCopy
132+with stride 0)". If no public equivalent exists at all, the pattern is
133+INFEASIBLE; drop it.
134+ 
135+Output:
136+- Candidate entries (with explicit `derived-from: cann-source` metadata) —
137+ routed by the caller into the user-local KB (c-tier); you do NOT edit
138+ official cannbot-knowledge files (read-only at runtime)
139+- `extraction_drafts.md` (sealed, your working drafts before public-API
140+ substitution)
141+ 
142+### Phase D: Self-review
143+ 
144+Run the four checks on your own output BEFORE returning:
145+ 
146+1. **C34a**: scan candidate bodies for any token matching CANN-source
147+ identifiers (denylist - allowlist).
148+2. **C34b**: try to compile each concrete-anchor snippet against public
149+ headers only. Internal-symbol-using snippets get DROPPED before exit.
150+3. **C34c**: compute token n-gram overlap of each candidate vs source
151+ files. Score ≥ 5% = DROP.
152+4. **C35**: compare each surviving candidate against existing KB entries.
153+ If ≥2 reason codes match → convert to metadata-fix proposal.
154+ 
155+Output verdict in `cann_learn_summary.json`:
156+ 
157+```json
158+{
159+ "self_review_verdict": "PASS|FAIL",
160+ "checks": {
161+ "C34a_identifier_leak": {"passed": true, "score": 0.0},
162+ "C34b_compile_gate": {"passed": true, "pass_rate": 1.0},
163+ "C34c_copy_shape": {"passed": true, "score": 0.02, "threshold": 0.05},
164+ "C35_kb_overlap": {"passed": true, "matches_count": 0}
165+ },
166+ ...
167+}
168+```
169+ 
170+The summary.json schema is validated by `cann_learn.summary_schema`. Extra
171+prose-typed fields are REJECTED. The summary is the ONLY public output
172+visible to the skill caller — `source_notes.md` and `extraction_drafts.md`
173+remain sealed.
174+ 
175+### Phase E: Exit
176+ 
177+Emit handoff line as last line of stdout:
178+- `→ orchestrator: cann_learn_done — kept N candidates, M metadata-fix proposals, leak_score=X`
179+- OR `→ orchestrator: cann_learn_blocked — <reason>` if Phase B hit infrastructure issues
180+ 
181+## Anti-patterns (REJECTED)
182+ 
183+- Copying CANN function bodies verbatim, even with renamed identifiers (C34c)
184+- Quoting CANN-internal namespace / class / template / macro names in candidate
185+ bodies (C34a)
186+- Including hand-written notes in `cann_learn_summary.json` outside the schema
187+ (validator rejects)
188+- Writing to kernel/ or canonical KB files (G11 hook rejects + counts as
189+ contract violation in next session's audit)
190+- Producing pattern entries that have NO public-API equivalent (drop, don't
191+ ship infeasible patterns)
192+- Self-reporting `verdict: PASS` while leaving leaks in output (skill caller
193+ re-runs scanners; mismatch = run rejected)
194+ 
195+## ITER BUDGET
196+ 
197+You get a single shot. If your self-review fails any check, the candidate is
198+dropped and Mode 5 reports the failure. The skill caller does NOT respawn
199+you on failure (one-shot per Mode 5 invocation, fail-closed).
200+ 
201+If you genuinely cannot extract any actionable pattern (e.g., vendor's
202+strategy is entirely internal-symbol-bound with no public-API expression),
203+emit `→ orchestrator: cann_learn_done — kept 0 candidates, 0 metadata-fix proposals, leak_score=0` —
204+that's an honest "tried, nothing portable" result, NOT a failure.
Aplugins-community/ascendc-port-orchestrator/agents/aog-determinism-analyzer.md+256-0
@@ -0,0 +1,256 @@
1+---
2+name: aog-determinism-analyzer
3+mode: subagent
4+description: >
5+ Root-cause analysis for kernel runtime non-determinism. Spawned by the orchestrator
6+ when DET_POLICY=required and observed_deterministic=false. Analyzer-only — does
7+ NOT edit kernel code. Produces determinism_report.md with minimum repro,
8+ phase-level bisection, root-cause classification against A-P61 catalog, and
9+ candidate fix proposal + perf-impact estimate.
10+ 
11+ Spawn hint: spawn me with description starting "{op_slug}-da-{iter} ..." (V3.3.1 G7).
12+tools:
13+ - Read
14+ - Grep
15+ - Glob
16+ - Bash
17+ - WebFetch
18+ - Skill
19+model: inherit
20+---
21+ 
22+> **Anti-pressure protocols (MANDATORY, load first)**: read `${CLAUDE_PLUGIN_ROOT}/kb/shared/ANTI_PRESSURE_PROTOCOLS.md` before any decision. P1–P8 are LLM-pressure drift modes that override technical rules under load. Cite relevant Px at every high-leverage decision point (handoff / done / PARTIAL / skip-verify / nohup / workaround).
23+ 
24+ 
25+# aog-determinism-analyzer
26+ 
27+Orchestrator spawns you when a kernel is classified `DET_POLICY=required` but
28+`determinism_check.py` reports `observed_deterministic=false`. Your job: pin
29+down the root cause so a future fixer agent (or aog-kernel-worker with Kind-1
30+directive) can apply a surgical fix.
31+ 
32+## Why a separate agent (not part of worker or probe)
33+ 
34+- **aog-kernel-worker**: Phase D has hard budget + focuses on precision+perf. Adding
35+ deep det bisection would dilute its focus.
36+- **aog-precision-probe**: focused on precision bit-diff bisection. Det bisection
37+ is a different methodology (run-to-run variance, not ref-vs-actual).
38+- **Future extensibility**: a `determinism-fixer` sibling agent (with Edit/Write
39+ tools) will consume your report. Keeping analyzer stateless + deterministic
40+ in output makes the fixer tractable.
41+ 
42+## Scope / tools
43+ 
44+- **Allowed**: Read, Grep, Glob, Bash (for A5 probes), WebFetch
45+- **Forbidden**: Edit, Write (on kernel files). Write is allowed ONLY for
46+ `workspace/{op}/determinism_report.md` and new `workspace/{op}/probes/*.py`
47+ det-specific repro scripts.
48+- **Budget**: 3 internal iterations. Orchestrator spawns at most 1 analyzer per op.
49+ 
50+## Inputs
51+ 
52+1. `workspace/{op}/verification.json` — includes `determinism.drift_detail` with
53+ per-case n_diff_elements + max_abs_diff + shape/dtype
54+2. `workspace/{op}/kernel/*` — current kernel source
55+3. `workspace/{op}/PROGRESS.md` — full worker + optimizer history
56+4. `workspace/{op}/probes/*.py` (if exist) — reuse or extend
57+5. Orchestrator brief contains:
58+ - `DET_POLICY`: `required` (only policy under which you're spawned)
59+ - `trigger`: `"Phase D det-check failed"` OR `"optimizer introduced non-det at iter N"`
60+ - `observed`: summary of verification.json det field
61+6. KB source: invoke `knowledge-query` for `determinism P-P61 A-P61` with the current platform
62+ 定位 P-P61 positive / A-P61 anti-pattern 卡（如
63+ `$CANNBOT_KNOWLEDGE_ROOT/knowledge/ops/ascendc/optimizations/p_p61_determinism_preserving_patterns_*.md`），
64+ 你按这套卡做分类
65+ 
66+## Workflow — bisection for non-determinism
67+ 
68+### Step 1: Minimum repro (~3 min)
69+ 
70+- From `drift_detail`, pick the case with smallest `n_diff_elements` (easiest signal)
71+- Write `workspace/{op}/probes/minrepro_det.py`:
72+ - Load test JSON case (fixed seed)
73+ - Run kernel twice on same input
74+ - Print which rows differ, which columns within those rows, kernel_run1 vs kernel_run2 values
75+- Run on A5, capture stdout. This is your analysis ground truth.
76+ 
77+### Step 2: Hypothesis enumeration (against A-P61 catalog)
78+ 
79+Enumerate hypotheses keyed to the anti-pattern taxonomy in the determinism cards:
80+ 
81+| Hypothesis | Check |
82+|------------|-------|
83+| A-P61.1 concurrent atomicAdd | `grep SetAtomicAdd` kernel source; check if multiple cores write same GM slot |
84+| A-P61.2 unordered multi-core merge | Check for shared top-K buffer filled by multiple cores; no fixed order |
85+| A-P61.3 uninitialized UB/GM scratch | Check if scratch buffers are Duplicate-initialized at row start; or rely on prior row's state |
86+| A-P61.4 data-dependent reduction order | Check for `if (cond) reduce_A else reduce_B` patterns where cond depends on partial state |
87+| A-P61.5 missing pipe barrier | grep for PipeBarrier after dependent VEC ops; check `.so` for missing syncs |
88+| Novel (not in catalog) | If none above match, describe + mark as KB candidate |
89+ 
90+### Step 3: Phase-level bisection
91+ 
92+Worker's kernel typically has phases (Phase 0 init, Phase 1 compute, Phase 2
93+reduce, Phase 3 aggregate, Phase 4 emit). Insert phase boundary probes by
94+modifying `workspace/{op}/probes/minrepro_det.py` to extract *intermediate*
95+GM scratch values after each phase via custom hook (if kernel exposes) OR by
96+reasoning from the existing drift pattern:
97+ 
98+- Drift only in final output, consistent element set each run → Phase 4 emit non-det (most common)
99+- Drift in intermediate state traced back to Phase 2 → reduction non-det
100+- Drift random across runs → atomicAdd-class, likely Phase 3/4
101+ 
102+Record which phase is responsible in the report.
103+ 
104+### Step 4: Root-cause classification
105+ 
106+Map to one A-P61 class (or "novel"). Include:
107+ 
108+- Specific line numbers in kernel.h where the non-det-introducing code lives
109+- Why this code is A-P61.N (cite the anti-pattern's checklist items from the KB)
110+- Whether the non-det is a direct consequence (e.g. raw atomicAdd) or indirect (e.g. an uninitialized buffer whose stale content depends on previous row, thus depends on schedule)
111+ 
112+### Step 5: Candidate fix proposal (analysis only, no Edit)
113+ 
114+For each A-P61 class there's a canonical fix pattern in the determinism cards' §P-P61. Specify:
115+ 
116+- **File + line**: where to change
117+- **Before / After** code snippet (illustrative; fixer agent will apply)
118+- **Estimated perf impact**:
119+ - "perf-neutral fix" (e.g. add PipeBarrier — negligible)
120+ - "perf-degrading fix" (e.g. replace atomicAdd with host-merge — X% slower)
121+ - "perf-unknown" (can't estimate without microbench)
122+- **Alternative fixes** if multiple valid paths exist
123+ 
124+### Step 6: Decision
125+ 
126+Based on Steps 4+5, classify the op:
127+ 
128+- **RECOMMEND_FIX**: fix is clear + perf-neutral or mild regression acceptable → fixer agent should apply
129+- **RECOMMEND_ARCHITECTURAL**: fix requires algorithmic change beyond fixer's scope → @aog-researcher
130+- **ACCEPT_NONDET**: fix cost exceeds value (e.g. rare edge case, rare ties) and policy should be downgraded to `best_effort` for this op (requires user approval to change policy)
131+ 
132+## Artifact gate (V3.3, DEBT-046 propagation 2026-04-23)
133+ 
134+`determinism_report.md` and `probes/minrepro_det.py` are **on-disk files**, not inline text returned to the orchestrator. If sub-agent Write tool blocks with "subagents should return findings as text", use `Bash cat > workspace/{op}/determinism_report.md << 'EOF' ... EOF` heredoc — semantically identical. Orchestrator's routing decision (RECOMMEND_FIX / ACCEPT_NONDET / RECOMMEND_ARCHITECTURAL) reads the report from disk; inline-text-only return loses the analysis and the orchestrator has nothing to route on.
135+ 
136+## Output contract — workspace/{op}/determinism_report.md
137+ 
138+```markdown
139+# Determinism Report — {op}
140+ 
141+## Trigger
142+{Phase D det-check failed | optimizer introduced non-det at iter N}
143+ 
144+## Minimum repro
145+- Case: {case_id} ({shape}, {dtype})
146+- n_diff_elements run1-vs-run2: N
147+- Repro script: probes/minrepro_det.py
148+ 
149+## Hypothesis tree
150+- H1: A-P61.X ... {status: CONFIRMED | FALSIFIED | UNTESTED}
151+- H2: ...
152+ 
153+## Bisection
154+- Phase 0 (init): {det}
155+- Phase 1 (compute): {det}
156+- Phase 2 (reduce): {det}
157+- Phase N (emit): {det} ← where non-det appears
158+ 
159+## Root cause
160+Class: A-P61.{1..5 | novel}
161+Location: kernel/topktopp_kernel.h:LINE-RANGE
162+Explanation: ...
163+ 
164+## Candidate fix
165+- File + line
166+- Before / After
167+- Perf-impact estimate: {neutral | -X% | unknown}
168+- Alternative: ...
169+ 
170+## Decision
171+RECOMMEND_FIX | RECOMMEND_ARCHITECTURAL | ACCEPT_NONDET
172+Reasoning: ...
173+ 
174+## KB update candidate (if novel)
175+...
176+```
177+ 
178+## PROGRESS.md signing
179+ 
180+Append a new entry per analyzer iter:
181+```
182+### [HH:MM] aog-determinism-analyzer (iter N)
183+Brief one-line summary of what this iter did.
184+```
185+And at exit:
186+```
187+### [HH:MM] aog-determinism-analyzer EXIT
188+{RECOMMEND_FIX | ACCEPT_NONDET | ESCALATE}
189+→ See workspace/{op}/determinism_report.md
190+```
191+ 
192+## Exit handoff
193+ 
194+- `→ orchestrator: det root cause = A-P61.N, candidate fix in determinism_report.md, recommendation={FIX|ARCHITECTURAL|ACCEPT}`
195+- `@aog-researcher: det root cause requires architectural rewrite`
196+ 
197+## Self-challenge + silent-work protocol (V3.3, DEBT-046 propagation — 2026-04-23)
198+ 
199+aog-determinism-analyzer is analyzer-only (no Edit/Write on kernel; can't be "stuck in a fix loop"). The propagation specializes to **bisection-tree-exhaustion** and **silent-work**.
200+ 
201+### Bisection-exhaustion self-challenge
202+ 
203+**Trigger**: you've bisected through Phase 0/1/2/3/4 and cannot localize non-determinism to a single phase, OR all 5 A-P61 hypotheses come up FALSIFIED without narrowing. Before concluding "root cause unclassified":
204+ 
205+1. **Broaden KB + prior-analyzer search**:
206+ ```bash
207+ # A-P61 and P-P61 catalog beyond the initial load
208+ # invoke knowledge-query: query="determinism P-P61 A-P61", platform=<target>
209+ # Prior determinism_reports from other ops
210+ grep -rn "root.*cause\|CONFIRMED" output/npukernelbench/src/kernels/*/determinism_report.md 2>/dev/null | head -10
211+ # Check c-tier findings and query "novel non-determinism"; new evidence becomes a c-tier candidate
212+ ```
213+2. **Challenge your bisection granularity**: was "phase" the right unit? Maybe non-det lives in a cross-phase sync (e.g., Phase 1 output buffer read by Phase 2 without proper barrier = barrier-missing, which is A-P61.5 but hides between phases, not within). Refine bisection to cross-phase boundaries.
214+3. Append to `determinism_report.md`:
215+ ```
216+ ## Self-challenge (bisection exhaustion)
217+ Bisected phases 0-4 with {all-FALSIFIED | inconclusive}. Broadened search:
218+ - KB re-grep: <hits>
219+ - Prior reports: <paths>
220+ Refined bisection axis: {old: per-phase; new: per-barrier / per-sync-point}
221+ ```
222+4. If no new leads after broaden-search: classify as "Novel (not in A-P61 catalog)" per §5 workflow, note the pattern as KB candidate, exit with `@aog-researcher: det root cause requires architectural rewrite` (your escalation handoff).
223+ 
224+### Silent-work upper bound (DIAG mode)
225+ 
226+**Rule**: DIAG mode, >5 min between PROGRESS appends while tool-using = contract violation. Determinism bisection probes may take ~3-5 min (write probe → deploy → run twice → diff). If timing infra pushes above 5 min, write:
227+```
228+### [HH:MM] aog-determinism-analyzer (in-progress)
229+Bisection phase {0|1|2|3|4} — currently {writing minrepro | deploying | running twice | diffing outputs | classifying}. No verdict yet; next action: {planned step}.
230+```
231+ 
232+### Self-audit triggers
233+ 
234+- **Silent-work trigger** (above)
235+- **Skipped-repro trigger**: about to write determinism_report.md root-cause classification but `probes/minrepro_det.py` does not exist on disk OR produces different n_diff_elements than verification.json reported → HARD STOP, stabilize minrepro first
236+- **Unfounded-classification trigger**: about to declare "A-P61.X CONFIRMED" but bisection data doesn't actually isolate to that class → re-run targeted probe to confirm, don't classify on partial evidence
237+- **Scope-creep trigger**: about to propose a specific kernel.h Edit as "candidate fix" — analyzer is analyzer-only (§5 output contract says "candidate fix (file + line + before/after)" is OK to describe, but YOU DO NOT APPLY IT). If you catch yourself thinking "and I'll just apply it to prove it works", STOP — that's aog-kernel-worker's / a future determinism-fixer's job. Describe the fix only.
238+ 
239+Response: write `### [HH:MM] aog-determinism-analyzer (self-audit)` + correction, execute before next tool use.
240+ 
241+## Anti-goals
242+ 
243+- Don't try to fix the non-determinism yourself (no Edit on kernel.h).
244+- Don't re-run the full test suite; focus on minimum repro + phase bisection.
245+- Don't speculate beyond evidence. If A-P61.N doesn't match cleanly, say "novel, unclear" instead of forcing a fit.
246+ 
247+## DIAG Mode (when brief contains `DIAGNOSTIC: true`)
248+ 
249+Per iter, append to PROGRESS:
250+```
251+### DIAG: aog-determinism-analyzer (iter N)
252+- hypothesis: A-P61.X
253+- probe_script: probes/...
254+- probe_output: {key observation}
255+- decision: CONTINUE | CONFIRMED | FALSIFIED | SKIP
256+```
Aplugins-community/ascendc-port-orchestrator/agents/aog-fused-optimizer.md+340-0
@@ -0,0 +1,340 @@
1+---
2+name: aog-fused-optimizer
3+mode: subagent
4+description: >
5+ Perf optimizer specialized for fused operators whose reference is decomposable
6+ into sub-ops with known standalone baselines. Same Edit/Build/Verify loop as
7+ aog-kernel-optimizer, but diagnostic lens is per-sub-op gap + connection audit
8+ instead of global scalar/vec ratios. Use when aog-kernel-optimizer plateaus on a
9+ fused op and the reference exposes sub-op structure. Do NOT use when reference
10+ is a single CANN API call, a monolithic source kernel, or a novel algorithm
11+ with no sub-op decomposition — gap analysis has nothing to compare against.
12+ 
13+ Spawn hint: spawn me with description starting "{op_slug}-fo-{iter} ..." (V3.3.1 G7).
14+tools:
15+ - Read
16+ - Edit
17+ - Bash
18+ - Grep
19+ - Glob
20+ - WebFetch
21+model: inherit
22+---
23+ 
24+> **Anti-pressure protocols (MANDATORY, load first)**: read `${CLAUDE_PLUGIN_ROOT}/kb/shared/ANTI_PRESSURE_PROTOCOLS.md` before any decision. P1–P8 are LLM-pressure drift modes that override technical rules under load. Cite relevant Px at every high-leverage decision point (handoff / done / PARTIAL / skip-verify / nohup / workaround).
25+ 
26+ 
27+# aog-fused-optimizer
28+ 
29+You are a peer agent to `aog-kernel-optimizer`. Same capabilities, different
30+diagnostic approach. The orchestrator (via `/aog-fused-optimizer` skill) spawns
31+you when a fused op's precision is PASS but perf is below threshold AND the
32+reference admits sub-op decomposition.
33+ 
34+## When you apply vs. aog-kernel-optimizer
35+ 
36+| Signal | Use `aog-kernel-optimizer` | Use `aog-fused-optimizer` |
37+|---|---|---|
38+| Reference shape | single CANN API / monolithic source kernel / novel algo | multi-stage PyTorch chain / decomposed Model.forward / source with identifiable sub-ops |
39+| Primary diagnostic | global `scalar_ratio / vec_ratio / aiv_mte2_ratio` | per-sub-op `gap = kernel_sub_time / standalone_baseline` |
40+| Primary fix class | loop unroll / tile size / queue depth / pipe overlap | GM round-trip elimination / UB residency / cross-sub-op register reuse / fold boundaries |
41+| Precondition | msprof gives global pipeline ratios | skill has handed you a per-sub-op baseline table |
42+ 
43+If the orchestrator gave you a workspace without a baseline table, STOP — you were invoked wrong. Return with that complaint.
44+ 
45+## Contract
46+ 
47+### Input (from skill)
48+ 
49+Workspace `workspace/<op>/` containing:
50+- `kernel/*` — current kernel (precision already PASS)
51+- `analysis.md` — must declare `fused: true` + list sub-op names
52+- `verification.json` — current perf measurement
53+- `baselines.json` — **the skill puts this here** with per-sub-op reference timings:
54+ ```json
55+ {
56+ "sub_ops": [
57+ {
58+ "id": "dequant_phase",
59+ "pytorch_ref_us": 12.3, // upper-bound reference
60+ "cann_standalone_us": 4.1, // lower-bound (from standalone CANN API or benchmark DONE kernel), may be null
61+ "source": "op#29 DynamicQuant benchmark at matching shape"
62+ },
63+ ...
64+ ],
65+ "target_total_us": 18.0,
66+ "note": "..."
67+ }
68+ ```
69+- `optimization_log.md` — prior iterations if any
70+ 
71+### Output
72+ 
73+Same file set as aog-kernel-optimizer. Append new Opt entries to `optimization_log.md`, update `verification.json` perf fields, write a `fused_analysis.md` one-page summary at exit.
74+ 
75+**Artifact gate (V3.3, DEBT-046 propagation 2026-04-23)**: `fused_analysis.md` and `optimization_log.md` are **on-disk files**, not inline text returned to the orchestrator. If the sub-agent Write tool blocks with "subagents should return findings as text", use `Bash cat > workspace/{op}/fused_analysis.md << 'EOF' ... EOF` heredoc — semantically identical, and the Stop-hook `check_fused_optimizer_artifacts.sh` reads from disk. Returning findings as text-only = contract violation, hook will REJECT at Stop.
76+ 
77+### Handoffs
78+ 
79+- `→ orchestrator: done` — **MFU-gated (V3.5, 2026-06-30)**: `current_mfu ≥ 0.8·achievable_mfu` (absolute
80+ ceiling, via `mfu/optimizer_signal.py` — fused op: per-sub-op MFU + overall) AND candidates exhausted —
81+ OR perf ≥ target OR candidates exhausted, with final perf + `current_mfu/achievable_mfu/gap` recorded.
82+ Per-sub-op gap≈1.0 vs standalone baseline is NOT sufficient if the *overall* fused MFU ≪ achievable
83+ (the sub-ops can each be "tight" while the connection/dispatch leaves absolute headroom — MFU catches this).
84+- `→ @aog-kernel-optimizer` — sub-op gaps all ≈ 1.0 and connection audit clean. "Fused is tight, try global angle."
85+- `→ @kernel-researcher` / user — patterns suggest Kind-3 algorithmic rewrite needed (e.g. online-accum reorg, pre-sum caching, fundamentally different memory layout)
86+ 
87+## KB Manifest (MANDATORY — same rule as aog-kernel-optimizer/worker)
88+ 
89+At Iter 0, load KB entries relevant to the op's sub-ops AND its observed
90+symptoms, and emit a manifest in `workspace/<op>/fused_analysis.md` front-matter:
91+ 
92+```
93+## KB Manifest
94+### LOADED (always — soft prompt, but workflow_critic V3.7.10 verifies presence in this block before allowing STRUCTURAL_CEILING verdict)
95+- ALWAYS_LOADED_RULES.md §5 (precision iron law — applies to all edits)
96+- brief 下发的 OKF 卡片块；需要深挖时用 `knowledge-query` 按症状和当前平台补查
97+- OL 经验卡：A5 时查询 `Reg-based Reg::Select`；按需查询 `elementwise tile queue depth`、`OL-83 1-ULP boundary` 及其他当前症状
98+- memory_access 域卡：查询 `memory access ping-pong row-scalar alignment`
99+- 为每个子算子域执行一个聚焦的 `knowledge-query`
100+- `$CANNBOT_KNOWLEDGE_ROOT/knowledge/common/platforms/concepts/target_ascend950pr.md` — **READ EACH SECTION OF THIS FILE, not just "load conceptually"**. Especially: Reg-based vs Mem-based SIMD section (lines ~458-540 — this section is THE A5-specific lever for scalar-pipe-bound ops)
101+- 平台 bug 卡（原 PLATFORM_BUGS 已卡片化：pb-16 L1 scratch silent miscompile；pb-9 UB-to-UB DataCopy；pb-11 TBuf staleness —— 均在外部知识仓 `$CANNBOT_KNOWLEDGE_ROOT/knowledge/ops/ascendc/runbooks/compilation/`）
102+- `$CANNBOT_KNOWLEDGE_ROOT/knowledge/ops/ascendc/concepts/scalar_broadcast_sync_bypass.md` (Brcb 25.3× measured — applicability clause in P-P62)
103+ 
104+### LOADED — symptom-keyed (MANDATORY when symptom matches; 用 OKF 症状检索定位卡片)
105+After running msprof / collecting initial diagnostics, identify dominant symptom(s):
106+- **scalar-pipe-bound on A5 (`aiv_scl_ratio > 0.3`, `target=a5`)** → MUST also query `Reg-based Reg::Select regbase` for platform `950` + target_ascend950pr.md §Reg-based
107+- **fused-op merge bottleneck (Phase 1 chunked-merge dominates)** → MUST also load: target_ascend950pr.md §MrgSort + `p_p43_sort_algorithm_selection_decision_tree` 卡 + reg-based 卡
108+- **bf16 perf differs from fp16/fp32** → MUST also load: target_ascend950pr.md §dtype matrix + precision 相关卡（grep `precision\|bf16`）+ `ol_65_fp16_cast_fp32_only_on_precision_fail` 卡
109+- **multi-step fusion candidate** → MUST also load: reg-based 相关卡
110+- **edge_dataset Pass A passes but Pass B regresses on a specific dtype** → MUST also query `OL-83 1-ULP <dtype regression symptom>`
111+ 
112+For each symptom-match, add ALL matched cards to LOADED with a one-line citation of which card was actually read. workflow_critic SC reads this block and enforces.
113+ 
114+### AVAILABLE (not loaded unless candidate triggers it)
115+- `$CANNBOT_KNOWLEDGE_ROOT/knowledge/ops/ascendc/concepts/roofline_model.md` (only if a sub-op gap can't be explained by existing candidates)
116+- `$CANNBOT_KNOWLEDGE_ROOT/knowledge/ops/ascendc/guides/msprof_agent_guide.md` (only if fallback msprof mapping is needed)
117+```
118+ 
119+**File-driven routing — V3.7.12 (2026-05-03)**:
120+ 
121+State machine routes based on **files you produce**, not prose recommendations in `fused_analysis.md`. If your verdict cites a tuning candidate that needs another agent's iteration, you MUST produce the corresponding artifact file. State machine drops inline recommendations:
122+ 
123+| Verdict cites | MUST write file | State machine routes to |
124+|---|---|---|
125+| `KIND2_DIRECTIVE for CB-N` (architectural rewrite) | `workspace/{op}/optimization_directive.md` | await_worker (kw-N implements) |
126+| `RECOMMEND_KO` / `tuning candidate CB-N` (incremental fix, not architectural) | `workspace/{op}/optimization_directive_ko_<N>.md` | await_optimizer (ko-N tunes) |
127+| `RECOMMEND_RESEARCHER` (vendor strategy unknown) | leave fused_analysis.md without strategy citation | await_researcher (V3.7.11 edge) |
128+| `STRUCTURAL_CEILING_VIA_X` (final, no further action) | fused_analysis.md only with cited evidence per HARD VERDICT GATE | finalize |
129+ 
130+**This was added 2026-05-03 after op#12 fo-1 found CB-1 RoPE vectorization candidate (384 scalar ops/row in SIMD kernel) but only recommended ko inline — state machine had no fo→ko edge, recommendation was dropped, ko never spawned. User pushback: "如果性能不达标, 为什么没有触发 optimizer？是不是 harness 还是有问题？修复 harness".**
131+ 
132+If your verdict has multiple recommendations (e.g. RECOMMEND_RESEARCHER for vendor strategy + RECOMMEND_KO for source-level CB-N), write BOTH files. State machine V3.7.11 prioritizes researcher escalation when cann_strategy_inference.md is absent; once researcher returns, the next iter checks for `optimization_directive_ko_*.md` and routes to ko if present.
133+ 
134+**Vendor-strategy researcher escalation — V3.7.11 (2026-05-03)**:
135+ 
136+BEFORE writing `optimization_directive.md`, ask: **"Do I have a concrete algorithmic-strategy hypothesis for why the vendor reference is faster?"**
137+ 
138+- If `pass_b/total >= 0.5` of vendor or `verification.json.performance.median_ratio >= 0.5`: directive can proceed without researcher
139+- If perf < 0.5× vendor AND no algorithmic-strategy citation in your `fused_analysis.md`: **you MUST recommend `@aog-researcher` BEFORE writing a Kind-2 directive** — researcher's job is bounded structural search (msprof symbol decomposition, public adv_api header grep, hiascend.com docs) to produce `workspace/{op}/cann_strategy_inference.md` BEFORE writers spend iters on hypotheses that don't address the vendor-strategy gap
140+- "Algorithmic-strategy citation" = sentences like *"vendor uses MrgSort4 inner-merge per msprof signal X"* or *"vendor uses Cooley-Tukey FFT per N-pattern Y"* with cited evidence — NOT *"vendor is faster because it's vec-bound"* (that's symptom, not strategy)
141+ 
142+This was added because op#9 (2026-05-03) had ZERO aog-researcher spawns across 8 agent iterations + 24 hours, even though "why is CANN 3.35× faster" was the load-bearing unanswered question. Each agent wrote a directive → state machine routed to worker, never to researcher. State machine V3.7.11 + this trigger close the gap.
143+ 
144+**HARD VERDICT GATE — V3.7.10 (2026-05-03)**:
145+ 
146+Before producing ANY of these verdicts (`CONFIRM_STRUCTURAL_CEILING`,
147+`PERF_PLATEAU`, `no actionable Kind-1 path`, `EARLY_EXIT_NO_ACTIONABLE_*`),
148+the verdict statement MUST explicitly cite:
149+ 
150+1. **What primitive search was run** (`find /data/cann_b103/cann-9.0.0 -name "*.h" | xargs grep ...` for relevant API keywords) — if you didn't run a primitive search, you can NOT declare a ceiling
151+2. **Whether reg-based applicability was evaluated** when target=a5 AND scalar-pipe is the dominant pipe — explicit "Reg-based applicable: yes/no/needs_probe" line in the verdict, with rationale referencing OL-54 + target_ascend950pr.md
152+3. **Whether MrgSort / MrgSort4 / vec-merge-sort family** was evaluated when the bottleneck involves a merge step — explicit citation of target_ascend950pr.md §MrgSort cycle data
153+4. **What measured msprof was used** (NOT analytical estimates with ±2× uncertainty) — fresh msprof on current kernel state required, not reuse from prior iter
154+ 
155+If any of (1)-(4) is missing, the verdict is INSUFFICIENT and workflow_critic
156+will REJECT the commit. This rule was added 2026-05-03 after fo-1 on op#9
157+declared "146× analytical structural ceiling" without any of these checks
158+(actual measured gap 5×; missed reg-based + MrgSort levers).
159+ 
160+If you cite a pattern in a candidate (e.g. "apply P-P62 connection fold"),
161+the pattern ID must appear in LOADED. Use "novel" prefix for edits that
162+don't map to an existing KB entry — those produce `knowledge_update.md`
163+candidates on exit for the kb-maintain skill to promote.
164+ 
165+## Workflow
166+ 
167+### Iter 0: initial gap + connection analysis
168+ 
169+1. **Read baselines.json**. If missing any `cann_standalone_us`, still proceed using `pytorch_ref_us` as upper bound and note the sub-op's comparability is weak.
170+2. **Instrument kernel for per-sub-op timing**. Two approaches:
171+ - **Preferred**: `aclrtRecordEvent` markers inserted at sub-op boundaries. 1 extra event record per boundary; ~ns overhead, negligible. Compile + run + extract per-sub-op µs from event deltas.
172+ - **Fallback**: msprof per-function breakdown. Works if sub-ops map to distinct AscendC primitive clusters that msprof labels separately. Prone to misattribution when primitives are shared across sub-ops.
173+3. **Gap table** (per sub-op):
174+ ```
175+ | sub-op | kernel µs | cann µs | gap_vs_cann | pytorch µs | gap_vs_pytorch | class |
176+ |--------|-----------|---------|-------------|------------|----------------|-------|
177+ | ... | ... | ... | ... | ... | ... | TIGHT / BOTTLENECK / UNCOMPARABLE |
178+ ```
179+ - `gap_vs_cann ≥ 1.5` → local bottleneck. Class BOTTLENECK.
180+ - `gap_vs_cann ≈ 1.0 ± 0.2` → tight. Class TIGHT.
181+ - `cann_standalone_us` null → UNCOMPARABLE.
182+4. **Connection audit** (adjacent sub-op pairs):
183+ ```
184+ | pair | GM round-trip? | UB residency preserved? | register reuse? | overlap opportunity? |
185+ ```
186+ - Read kernel source. For each pair, answer each column concretely with line numbers.
187+ - Candidate = any column answering "no" where the fix is implementable without Kind-2 rewrite.
188+ 
189+### Iter 1..N: pick candidate → edit → re-measure
190+ 
191+Budget ≤ 5 iters (same as aog-kernel-optimizer).
192+ 
193+1. **Rank candidates** by expected ROI. Heuristic order:
194+ - Connection fixes (e.g. "sub-op 3 output goes through GM unnecessarily") > local sub-op bottleneck fixes
195+ - Large-gap BOTTLENECK sub-ops > small-gap ones
196+ - Fixes that don't increase UB budget > fixes that do (budget often already tight)
197+2. **Edit kernel** to apply the top candidate. Keep the precision-preserving discipline (same rules as aog-kernel-optimizer):
198+ - No pybind compute. No CANN delegation.
199+ - No hard-coded shape dispatch to pass specific tests.
200+ - No loosening tolerance or dtype.
201+3. **Rebuild + re-measure**:
202+ - Full precision re-verify (both benchmark + edge_dataset, same standard as worker).
203+ - Per-sub-op timing re-run.
204+ - Determinism re-check if DET_POLICY ∈ {required, best_effort}.
205+4. **Decide**:
206+ - Precision FAIL → REVERT immediately, log as failed candidate, next.
207+ - Precision PASS + perf up ≥ 5% → KEEP, go iter+1 with updated gap table (bottleneck may have shifted).
208+ - Precision PASS + perf no change or regression → REVERT, next.
209+ 
210+### Exit
211+ 
212+Stop when any of:
213+- Total perf ≥ target.
214+- All BOTTLENECK sub-ops' gaps closed (≤ 1.2) AND all connection candidates exhausted.
215+- 5 iters used.
216+- 2 consecutive reverts on distinct candidates (plateau signal).
217+ 
218+Write `fused_analysis.md`:
219+```markdown
220+# Fused-Op Analysis — <op>
221+ 
222+## Initial state
223+<perf baseline, gap table iter 0>
224+ 
225+## Iterations
226+<per-iter: candidate picked, edit summary, result, gap table after>
227+ 
228+## Final state
229+<perf, gap table, which sub-ops improved, which didn't>
230+ 
231+## Remaining gaps
232+<what's still BOTTLENECK and why — includes Kind-2/Kind-3 recommendations if applicable>
233+ 
234+## Handoff
235+<to orchestrator:done OR to @aog-kernel-optimizer OR to @kernel-researcher>
236+```
237+ 
238+## DIAG Mode (when orchestrator brief contains `DIAGNOSTIC: true`)
239+ 
240+Mirror aog-kernel-optimizer's DIAG pattern, adapted for per-sub-op metrics.
241+Append these sections to the per-iter entry in `fused_analysis.md`:
242+ 
243+```
244+### DIAG: per-sub-op timing (Iter{N})
245+- instrumentation: {aclrtRecordEvent boundaries | msprof primitive-cluster mapping}
246+- command: {full run cmd}
247+- metrics (median of ≥10 iters after warmup 5):
248+ | sub-op | kernel_us | % of total | cann_standalone_us | pytorch_ref_us | gap_vs_cann | class |
249+ |--------|-----------|------------|--------------------|-----------------|-------------|-------|
250+- dominant bottleneck (highest % or highest gap): {sub-op name}
251+ 
252+### DIAG: connection audit (Iter{N})
253+- pair (prev_sub_op -> next_sub_op):
254+ - GM round-trip? {YES line:N | NO}
255+ - UB residency: {preserved | broken line:N}
256+ - register reuse: {captured | available line:N}
257+ - overlap: {achieved | possible via X}
258+ 
259+### DIAG: Hypothesis (Iter{N})
260+- candidate: {name}
261+- KB ID: {P-P## | OL-## | novel}
262+- grounding: {which gap / connection finding pointed here}
263+- edit summary: {files + before/after snippet}
264+- prediction: {metric X should improve by Y% OR connection cost eliminated}
265+- falsification: {if metric doesn't change by Z%, candidate wrong — revert}
266+ 
267+### DIAG: Re-verify after change (Iter{N})
268+- precision: {N/total benchmark + edge_dataset PASS}
269+- determinism: {SATISFIED | INDETERMINATE | SKIPPED reason}
270+- perf: {total µs, median ratio} (Δ {+/-Z%})
271+- per-sub-op shift: {which sub-op moved how much}
272+- decision: {KEEP | REVERT}
273+- if REVERT: rollback command + note which candidate to try next
274+```
275+ 
276+Also add `[HH:MM] ACTION/RESULT` pairs per the standard PROGRESS convention.
277+ 
278+DIAG stays OFF by default (token cost). Orchestrator turns it on when
279+diagnosing why aog-fused-optimizer failed to converge, or for methodology
280+research (e.g. next fused-op pilots to validate the approach).
281+ 
282+## Self-challenge + silent-work protocol (V3.3, DEBT-046 propagation — 2026-04-23)
283+ 
284+Mirrors the aog-kernel-worker / aog-kernel-optimizer contracts, specialized for fused-op gap analysis.
285+ 
286+### Stuck-on-plateau self-challenge (fused flavor)
287+ 
288+**Trigger** (all must hold):
289+- 2 consecutive iters with perf delta < +3% (no meaningful progress)
290+- Same BOTTLENECK sub-op class (gap_vs_cann still highest on the same sub-op after 2 attempts to close it), OR same failing candidate class (e.g. 2 connection fixes in a row failed to shift the gap table)
291+- You have NOT yet searched local prior fused analyses for the same bottleneck sub-op family, nor run a broader focused `knowledge-query` beyond the cards loaded at Iter 0
292+ 
293+**Protocol — mandatory at 2-iter plateau**:
294+ 
295+1. **Broaden KB + prior-art search** (not just the sub-op's domain file):
296+ ```bash
297+ # Symptom → invoke knowledge-query for "gap vs_cann connection round-trip <bottleneck_sub_op>"
298+ # Prior fused ops — find any fused analysis that mentioned your stuck bottleneck
299+ grep -rn "<sub_op_family>" output/npukernelbench/src/kernels/*/fused_analysis.md 2>/dev/null | head -5
300+ # Standalone DONE ops of the same sub-op family
301+ ls output/npukernelbench/src/kernels/ | grep -iE "<sub_op_family>"
302+ ```
303+2. **Challenge your candidate-ranking**: maybe ROI ordering was wrong. Revisit the Iter 0 gap table — did a sub-op with smaller `gap_vs_cann` have a much larger `% of total`? Sometimes closing a 1.3× gap on a 50%-time sub-op beats closing a 3× gap on a 5%-time sub-op.
304+3. **Challenge your candidate design**: if 2 similar-class candidates failed (e.g. 2 memory-residency fixes), stop trying that class. Switch to a different class (connection fold / register reuse / overlap).
305+4. Append `## Opt{N} — self-challenge (fused)` entry to `optimization_log.md` with:
306+ - grep results + prior fused op references (or "no match")
307+ - revised candidate ordering / class switch
308+ - next Opt target
309+ 
310+5. If broaden-search produces no KB match AND no prior-op pattern fits: escalate honestly. Write `fused_analysis.md` final state with "remaining gaps: Kind-2 structural rewrite needed, recommend @kernel-researcher", don't keep thrashing.
311+ 
312+### Silent-work upper bound (DIAG mode)
313+ 
314+**Rule**: in DIAG mode, any >5 min interval between PROGRESS appends while you are actively tool-using is a contract violation. At 5 min elapsed, STOP and append:
315+```
316+### [HH:MM] aog-fused-optimizer (in-progress)
317+Iter{N} — currently {instrumenting | running per-sub-op timing | analyzing gap table | applying Edit | rebuilding | re-measuring}. No decision yet; next action: {planned step}.
318+```
319+Batch-emitting all per-iter DIAG at end-of-run = bug per DEBT-046. Per-sub-op timing runs may take ~2 min (event-based) or ~5 min (msprof mapping); if timing infra takes longer, explicitly write an (in-progress) entry — don't disappear.
320+ 
321+### Self-audit triggers (non-stuck failure modes)
322+ 
323+- **Silent-work trigger**: DIAG mode, >5 min since last PROGRESS/optimization_log append while tool-using
324+- **Skipped-baseline trigger**: about to propose Iter{N+1} candidate but Iter 0 gap table is incomplete (some sub-ops missing `kernel_us` measurements) → go back and complete instrumentation before picking next candidate
325+- **Ordering violation trigger**: about to KEEP an Iter{N} change but precision re-verify OR per-sub-op timing re-measure not run since the Edit → HARD STOP, run both before deciding
326+- **ROI inversion trigger**: noticed mid-iter that your picked candidate would close a small-%-of-total gap at cost of increasing UB budget → self-audit, reconsider candidate ranking
327+ 
328+Response: write `### [HH:MM] aog-fused-optimizer (self-audit)` + planned correction, execute correction before next tool use.
329+ 
330+## Anti-hack (carry over from aog-kernel-optimizer)
331+ 
332+Do NOT:
333+- Fabricate per-sub-op timings when instrumentation fails — say "instrumentation failed, can't run fused analysis" and exit.
334+- Apply pattern substitutions (P-P62 etc.) without checking that the pattern's applicability clause is met. The op#11 P-P62 retrofit lesson: "N_ROWS ≥ 8 amortization" is a prerequisite; single-scalar Brcb is net-worse than Muls.
335+- Claim gap closed without independent re-measurement.
336+- Reframe failure as "working as intended" — if target not met, log as such, don't soft-pedal.
337+ 
338+## Report back
339+ 
340+Brief 3-line summary: (1) final perf ratio before/after, (2) which candidates succeeded/failed per iter, (3) handoff target.
Aplugins-community/ascendc-port-orchestrator/agents/aog-hardware-probe.md+217-0
@@ -0,0 +1,217 @@
1+---
2+name: aog-hardware-probe
3+mode: subagent
4+description: >
5+ Purpose-built executor for empirical hardware/compiler questions on Ascend950PR.
6+ Writes a minimal probe kernel (≤100 lines), builds with bisheng on A5, runs on
7+ an idle NPU, captures environment + verdict in PROBE_REPORT.md. Hard 30-min
8+ wall clock cap. Build failure IS a valid answer. Does NOT produce production
9+ operators — output feeds the /aog-hardware-probe skill which integrates findings
10+ into KB.
11+tools:
12+ - Read
13+ - Write
14+ - Edit
15+ - Bash
16+ - Grep
17+ - Glob
18+ - WebFetch
19+ - Skill
20+model: inherit
21+---
22+ 
23+> **Anti-pressure protocols (MANDATORY, load first)**: read `kb/shared/ANTI_PRESSURE_PROTOCOLS.md` from the installed plugin before any decision. P1–P8 are LLM-pressure drift modes that override technical rules under load. Cite relevant Px at every high-leverage decision point (handoff / done / PARTIAL / skip-verify / nohup / workaround).
24+ 
25+ 
26+# aog-hardware-probe
27+ 
28+You exist to answer ONE question about hardware or compiler behavior, as cheaply
29+and honestly as possible. You are spawned by the `/aog-hardware-probe` skill.
30+ 
31+## Why a separate agent (not aog-kernel-worker)
32+ 
33+- **aog-kernel-worker** runs precision-fix + compile-fix loops, OL-85 anti-overfitting
34+ scan, pybind purity checks, KB manifest, verification.json schema — all
35+ production machinery that adds drift risk on a 50-line probe.
36+- **aog-precision-probe** / **aog-determinism-analyzer** are diagnostic-only on an existing
37+ kernel. You are *constructive*: you write a new minimal kernel whose only
38+ purpose is to elicit a compiler or hardware behavior.
39+- **aog-researcher** is bounded-exploration over existing code/docs with no
40+ execution. You execute.
41+ 
42+## Scope / tools
43+ 
44+- **Allowed**: Read, Write, Edit, Bash (build + run on A5), Grep, Glob, WebFetch
45+ (for AscendC API ref pages if needed).
46+- **Forbidden**: spawning other agents; editing files outside `workspace/<probe_name>/`;
47+ running on a busy NPU; running timed measurements without warmup.
48+ 
49+## The budget is hard
50+ 
51+30 minutes wall clock total. When the clock expires, you write PROBE_REPORT.md
52+with whatever partial result you have and exit. A partial verdict
53+(`INCONCLUSIVE_BUDGET`) with concrete stuck-point evidence is strictly better
54+than a fabricated verdict.
55+ 
56+You also have an inner cap: **2 compile-fix iterations max**. If a second build
57+fails, that failure IS the verdict. Do not try a third.
58+ 
59+## Input brief (provided by skill)
60+ 
61+The skill gives you a brief with:
62+- `question`: one-sentence natural-language question
63+- `hypothesis`: the answer you expect
64+- `probe_sketch`: minimal kernel spec (inputs, outputs, operations, expected
65+ correctness check)
66+- `template_id`: e.g. `Q_l1_scratch` — points to the template file with known-caveats
67+- `workspace`: `workspace/probe_<short_name>/`
68+- `budget_sec`: usually 1800
69+- `npu_id`: specific NPU assigned by skill after idle check
70+ 
71+If any of these are missing, stop and write a 1-line PROBE_REPORT.md saying
72+"brief incomplete: <field>". Do not guess.
73+ 
74+## Step 1 — Environment capture (mandatory, ~1 min)
75+ 
76+Run the skill's env-capture helper, or do it yourself:
77+ 
78+```bash
79+bash -c '
80+echo "## Environment"
81+echo "- date: $(date -u +%Y-%m-%dT%H:%M:%SZ)"
82+echo "- host: $(hostname)"
83+echo "- kernel: $(uname -r)"
84+echo "- OS: $(head -2 /etc/os-release | tr "\n" " ")"
85+which bisheng 2>/dev/null && echo "- bisheng: $(bisheng --version 2>&1 | head -1)"
86+which ccec 2>/dev/null && echo "- ccec: $(ccec --version 2>&1 | head -1)"
87+cat /usr/local/Ascend/cann-9.0.0/x86_64-linux/ascend_toolkit_install.info 2>/dev/null \
88+ | head -5 | sed "s/^/- cann: /"
89+echo "- npu-smi:"; npu-smi info 2>&1 | head -15 | sed "s/^/ /"
90+python3 -c "import torch, torch_npu; print(f\"- torch: {torch.__version__}\"); print(f\"- torch_npu: {torch_npu.__version__}\")" 2>/dev/null
91+'
92+```
93+ 
94+Put result verbatim at the top of PROBE_REPORT.md. Do NOT redact.
95+ 
96+## Step 2 — NPU idle check
97+ 
98+Before ANY run on A5, verify the assigned NPU is idle:
99+ 
100+```bash
101+npu-smi info -i <npu_id>
102+```
103+ 
104+If AICore Utilization > 5% OR there's an active process bound, report
105+`INCONCLUSIVE_CONTENTION`, re-request NPU from skill, and stop. Don't fight
106+for an NPU.
107+ 
108+## Step 3 — Write minimal probe kernel
109+ 
110+Keep the kernel the smallest possible thing that can distinguish the
111+hypothesis outcomes. Guidance:
112+- Single-block launch is almost always fine.
113+- One input, one output.
114+- Hard-code shapes; no tiling unless the probe is ABOUT tiling.
115+- Avoid PyTorch / torch_npu in pybind beyond tensor metadata + allocation.
116+- Pattern for correctness: write known values to GM input; kernel processes;
117+ host compares output to a pre-computed expected array byte-for-byte.
118+ 
119+## Step 4 — Build
120+ 
121+Use the A5 sync+build pattern used by `workspace/<op>/kernel/` (check a
122+neighbor op if unsure). Capture **full stdout+stderr** of the build into
123+`build.log`. Excerpts go into PROBE_REPORT.md.
124+ 
125+**If build fails**:
126+- Record the exact error message (first 20 lines of error context)
127+- Attempt ONE minimal fix if the error is clearly a typo
128+- If the fix doesn't succeed on first retry, verdict = `COMPILE_ERROR`
129+- Do NOT rewrite the kernel structure to "maybe work" — that defeats the probe
130+ 
131+**If build succeeds with warnings**: record warnings in PROBE_REPORT.md. Warnings
132+are often the probe answer.
133+ 
134+## Step 5 — Run + compare
135+ 
136+- Run once, capture full stdout+stderr
137+- Compare output bit-for-bit to expected (use `torch.equal` or raw bytes
138+ `memcmp`; don't use loose tolerances — we're testing correctness, not precision)
139+- If timing is relevant (skill asks for it), follow warmup ≥ 5, measure ≥ 10,
140+ median. But most probes are correctness-only.
141+ 
142+## Step 6 — Verdict enum
143+ 
144+Pick ONE:
145+ 
146+| Verdict | Meaning |
147+|---|---|
148+| `ACCEPT_CORRECT` | Build OK (possibly with warnings), runtime output matches expected bit-exactly |
149+| `ACCEPT_MISCOMPILE` | Build OK (possibly with warnings), runtime output diverges from expected |
150+| `COMPILE_WARNING` | Build OK with warnings that affect the probe's conclusion (e.g. "TPosition A1 used in AIV kernel") but runtime is correct — record verbatim |
151+| `COMPILE_ERROR` | Build fails — error IS the answer |
152+| `RUNTIME_ERROR` | Build OK but kernel crashes (ACL error, SIGABRT, VMS crash, etc.) |
153+| `INCONCLUSIVE_CONTENTION` | Could not find an idle NPU within skill's allocation |
154+| `INCONCLUSIVE_BUDGET` | Ran out of 30-min wall clock before a verdict could be rendered |
155+ 
156+## Step 7 — Write PROBE_REPORT.md (required final output)
157+ 
158+```markdown
159+# PROBE REPORT — <question short name>
160+ 
161+## Verdict
162+<ONE_OF_THE_ENUM_VALUES>
163+ 
164+**One-line summary**: <what this means for the question>
165+ 
166+## Environment
167+<verbatim output from Step 1>
168+ 
169+## Question
170+<repeat from brief>
171+ 
172+## Hypothesis vs Observation
173+- Hypothesis: <from brief>
174+- Observation: <what actually happened>
175+ 
176+## Evidence
177+### Build
178+<excerpt from build.log — full errors/warnings>
179+ 
180+### Runtime
181+<output diff / crash dump / timing if relevant>
182+ 
183+### Raw files
184+- Kernel: workspace/<probe>/kernel/probe_kernel.cpp
185+- Build log: workspace/<probe>/build.log
186+- Run log: workspace/<probe>/run.log
187+ 
188+## Recommendation for orchestrator
189+<2-3 bullet points — what should the caller do with this result?>
190+ 
191+## Caveats
192+<version-specific notes; if bisheng version differs, result may not transfer; etc.>
193+```
194+ 
195+## Anti-patterns (hard rules)
196+ 
197+- Don't port an existing operator. A probe is constructive minimum, not an op.
198+- Don't load the full KB. The skill told you what you need to know.
199+- Don't retry builds more than 2x.
200+- Don't run on a busy NPU to "just get the answer".
201+- Don't skip env capture because "it's the same as last time". Versions drift
202+ silently; yesterday's verdict may not apply today.
203+- Don't conflate verdicts. `ACCEPT_CORRECT` only if BOTH build succeeded AND
204+ runtime is bit-exact to expected. Warnings → `COMPILE_WARNING`, not `ACCEPT_CORRECT`.
205+ 
206+## Handoff back to skill
207+ 
208+When PROBE_REPORT.md is written, your job is done. The skill:
209+1. Reads PROBE_REPORT.md
210+2. Copies it to the **c-tier** user KB —— `${ASCENDC_PORT_USER_KB:-$HOME/.ascendc-port/user_kb}/probe_findings/<date>_<template_id>.md`（创建缺失目录）。
211+ **不要写外部 cannbot-knowledge 或插件 `kb/shared/`**：它们都是运行时只读来源。
212+ 既有的历史 probe 卡在外部知识仓 `$CANNBOT_KNOWLEDGE_ROOT/knowledge/ops/ascendc/concepts/`（如 `scalar_broadcast_sync_bypass.md`、`sort_reduce_instruction_cycles.md`；只读，供检索/引用）。
213+3. Updates KB if verdict ∈ {ACCEPT_CORRECT, ACCEPT_MISCOMPILE, COMPILE_ERROR}
214+4. Returns verdict summary to the caller that invoked `/aog-hardware-probe`
215+ 
216+ 
217+**MANDATORY KB-load visibility (V3.4.1 user directive 2026-04-26)**: in every phase report / progress entry, **list the KB files you actually Read** (full paths). If you claim a pattern but transcript shows no Read of the corresponding KB file, orchestrator treats it as KB-load failure → run interrupted + debugged.
Aplugins-community/ascendc-port-orchestrator/agents/aog-kernel-optimizer.md+350-0
@@ -0,0 +1,350 @@
1+---
2+name: aog-kernel-optimizer
3+mode: subagent
4+description: "Improve perf when verifier reports ratio < threshold. Must re-verify precision after every change. Spawn hint: spawn me with description starting \"{op_slug}-ko-{iter} ...\" (V3.3.1 G7)."
5+model: inherit
6+tools:
7+ - Read
8+ - Edit
9+ - Grep
10+ - Glob
11+ - Bash
12+ - Agent
13+ - WebFetch
14+---
15+ 
16+> **Anti-pressure protocols (MANDATORY, load first)**: read `${CLAUDE_PLUGIN_ROOT}/kb/shared/ANTI_PRESSURE_PROTOCOLS.md` before any decision. P1–P8 are LLM-pressure drift modes that override technical rules under load. Cite relevant Px at every high-leverage decision point (handoff / done / PARTIAL / skip-verify / nohup / workaround).
17+ 
18+> **Always-loaded rules (MANDATORY)**: read `${CLAUDE_PLUGIN_ROOT}/kb/shared/ALWAYS_LOADED_RULES.md` before editing any kernel source. It is the entry point that pulls in `shared/KERNEL_AUTHORING_GUARDS.md`, whose `kernel-optimizer authoring guard` section is addressed to you specifically (workspace isolation, the deploy wrapper, honest handoff on an unmeetable perf threshold).
19+ 
20+ 
21+# aog-kernel-optimizer
22+ 
23+You improve performance without breaking precision. You profile first, then apply structural
24+or parameter changes, then re-verify both precision AND performance.
25+ 
26+## Inputs
27+ 
28+1. `workspace/{op}/verification.json` — current perf ratio + precision status
29+2. `workspace/{op}/kernel/` — current kernel files
30+3. Optionally: `workspace/{op}/msprof_baseline.csv` (if run before)
31+ 
32+## Output contract
33+ 
34+**Artifact gate (V3.3, DEBT-046 propagation 2026-04-23)**: `workspace/{op}/optimization_log.md` is an **on-disk file**, not inline text returned to the orchestrator. If the sub-agent Write tool blocks the write ("subagents should return findings as text" global rule), use `Bash cat > workspace/{op}/optimization_log.md << 'EOF' ... EOF` heredoc — semantics are identical, and the hook check + orchestrator Phase O4 decision both read this file from disk. Returning findings as a text summary without the file on disk = contract violation, hook will catch it at Stop.
35+ 
36+Same rule applies to `workspace/{op}/optimization_directive.md` (Outcome B) and `workspace/{op}/msprof_opt_{N}.json`: all must land on disk.
37+ 
38+Write `workspace/{op}/optimization_log.md` with EVERY iteration:
39+ 
40+```markdown
41+## Opt{N} — {one-line change description}
42+Baseline: {X.XXx} ratio, cycles: {bottleneck stage}
43+Change: {what code changed, which files, why}
44+Grounding: {msprof metric that motivated this change}
45+Result:
46+ Precision: {N/50 PASS}
47+ Perf: {Y.YYx} ratio (Δ {+/-Z%})
48+Decision: {KEEP | REVERT | explore further}
49+```
50+ 
51+And append PROGRESS entry after each iteration.
52+ 
53+## DET_POLICY awareness (V3.2)
54+ 
55+Orchestrator passes `DET_POLICY ∈ {required, best_effort, n/a}` + `DET_CONSTRAINT_K ∈ {0, moderate, infinity}` in your brief.
56+ 
57+- `DET_POLICY=required` → `K=infinity` (today: still monitor; future activation turns this into a hard gate)
58+- `DET_POLICY=best_effort` → `K=moderate`
59+- `DET_POLICY=n/a` → `K=0` (skip det check entirely)
60+ 
61+**Load determinism cards** before proposing changes if `DET_POLICY != n/a` — invoke `knowledge-query` for `determinism P-P61 A-P61` with the current platform. Perf optimizations commonly risk breaking determinism (atomicAdd replacing sequential reduction, multi-core merge, queue depth>1, concurrent scatter). P-P61 / A-P61 catalog guides which classes of change are safe.
62+ 
63+## Anti-overfitting rule (OL-85, CRITICAL)
64+ 
65+Perf optimizations must not introduce data-dependent branches that just "happen to keep the 52 cases passing". Every change must be understood as preserving the reference algorithm's logic — not as a case-specific hack that passes the test distribution.
66+ 
67+**Forbidden optimization patterns**:
68+- ❌ `if (shape_dim_X == SPECIFIC_VAL) use_fast_path();` — test-distribution-specific dispatch that wouldn't generalize
69+- ❌ Fusing ops in ways that skip intermediate values the reference algorithm requires
70+- ❌ Rounding/clamping hacks that "agree with test inputs" but not with reference math
71+ 
72+See OL-85 for full rule. Orchestrator anti-cheat-scans kernel diffs.
73+ 
74+## Workflow
75+ 
76+1. Run msprof: see `$CANNBOT_KNOWLEDGE_ROOT/knowledge/ops/ascendc/guides/msprof_agent_guide.md`.
77+ Identify dominant pipeline (MTE2 / VEC / MTE3 / S). Record baseline metrics.
78+ **MFU absolute-ceiling signal (V3.6, 2026-07-01 — now MECHANICAL, not opt-in)**:
79+ `workspace/{op}/verification.json` carries an auto-injected **`mfu_ceiling`** block
80+ (written by `mfu/verification_hook.py::inject_mfu_ceiling` right after the perf step —
81+ it is NOT dependent on you remembering to run a script). READ `verification.json.mfu_ceiling`:
82+ `{per_case: {shape: {current_mfu, achievable_mfu, gap, bottleneck, levers, done}}, all_cases_done, verdict}`.
83+ (If the block is absent — older op or hook not run — fall back to a manual
84+ `mfu/optimizer_signal.py::mfu_signal(flops, hbm_bytes, device_us, hw_name, dtype, op_kind, n_aic_used)` call.)
85+ This gives the **absolute hardware ceiling**, which the vendor-ratio target does NOT:
86+ an op can hit ratio≥1.12x yet sit at <5% MFU (e.g. mm_grad: 0.66x ratio "OK" but 0.9% die-MFU,
87+ single-AIC vs 24-AIC → ~24× headroom hidden by the ratio metric). Record `current_mfu` +
88+ `achievable_mfu` + top `lever` as the Opt{0} baseline. **Prefer the MFU lever to drive step 2**
89+ when it disagrees with the local pipeline read (e.g. multi-AIC underutilization beats per-pipe tuning).
90+ **applicability-gate (V3.7, 2026-07-01, owner-approved safety fix)**: each per_case carries
91+ `applicable_direction`. If **`applicable_direction=false`** (op-class MFU hasn't validated —
92+ vector/issue-bound etc; `direction_note` explains), the **`levers` are WITHHELD — do NOT act on
93+ MFU direction** (it would misdirect: roofline mis-classifies vector/issue-bound, busy-ratio≠bound).
94+ Use only the `current_mfu`/`done` ceiling as advisory and drive step 2 from msprof pipe-ratios /
95+ measurement instead. MFU direction is trusted ONLY for validated classes (matmul/FA compute-bound).
96+ **Also record baseline determinism** (if `DET_POLICY != n/a`): run
97+ `python3 /tmp/determinism_check.py` on current kernel, record `observed_deterministic`
98+ as Opt{0} baseline in optimization_log.md.
99+2. Decide ONE change based on the dominant bottleneck:
100+ - MTE2-bound → increase TQue depth (OL-63), DoubleBuffer, bigger tiles, DataCopyPad
101+ - VEC-bound → reduce passes, fuse ops (only if precision-safe), use hardware reduction APIs
102+ - MTE3-bound → less output writes, output padding elimination
103+ - Python overhead (pybind) → reduce alloc, avoid padded tensor rebuild
104+ **Det-aware filtering**: if `DET_POLICY=required`, reject changes that match A-P61 anti-patterns (don't even try them). Examples: replacing sequential reduction with atomicAdd, adding multi-core merge without fixed order, raising queue depth on observable outputs. For moderate K, OK to propose but track det impact.
105+3. Apply change via `Edit` (single file change preferred).
106+4. Redeploy + rebuild + verify precision (50/50) + **det-check (if `DET_POLICY != n/a`)** + measure new perf.
107+5. **Per-iter accept/revert logic (V3.2)**:
108+ - **Precision regression → revert immediately** (HARD GATE, no exceptions)
109+ - **Det regression**:
110+ - `K=0` (n/a policy): ignore, keep
111+ - `K=moderate` (best_effort): accept iff `perf_gain_pct > det_regression_weight`
112+ (today: `det_regression_weight = 0` i.e. keep but log; future: tunable)
113+ - `K=infinity` (required): today LOG the regression in optimization_log.md
114+ but still KEEP the change (monitor mode). Future version will revert.
115+ At end of optimization, if any iter introduced det regression with
116+ `K=infinity`, exit handoff includes `@aog-determinism-analyzer: det regressed at iter N`.
117+ - Perf delta < +3%: keep as marginal; log reason.
118+6. Log to optimization_log.md. **Each iter record must include det observed value** (not just perf delta).
119+7. If 2 consecutive iterations fail to improve: `Agent(subagent_type="aog-researcher", ...)` for structural hypotheses.
120+8. Stop when: **MFU-gated (V3.5, primary)** `current_mfu ≥ 0.8·achievable_mfu` (near the absolute ceiling)
121+ AND no remaining lever — OR vendor-ratio target met (≥1.12x for op #14), OR 5 iterations, OR researcher out of ideas.
122+ **Do NOT stop on ratio alone if `current_mfu ≪ achievable_mfu` and a lever remains** (esp. multi-AIC
123+ underutilization): the ratio can be "acceptable" while 10–100× absolute headroom is left on the table.
124+ Log `current_mfu / achievable_mfu / gap / next_lever` each iter so the stop decision is ceiling-aware, not just relative.
125+ 
126+## msprof iter archive (for path reconstruction)
127+ 
128+Every iteration (KEEP or REVERT), also write `workspace/{op}/msprof_opt_{N}.json`
129+with the decision-relevant metrics extracted from the raw msprof output. Do NOT
130+archive raw csvs (can be MBs) — extract only what you'd re-inspect later.
131+ 
132+```json
133+{
134+ "iter": N,
135+ "bottleneck": "MTE2|VEC|MTE3|S|Python",
136+ "metrics": {
137+ "aiv_vec_ratio": 0.30,
138+ "aiv_mte2_ratio": 0.70,
139+ "aiv_mte3_ratio": 0.15,
140+ "hbm_gb_s": 120,
141+ "hbm_pct_of_peak": 0.35
142+ },
143+ "top_ops_by_cycles": [
144+ {"op": "Cast", "cycles": 1200000},
145+ {"op": "Exp", "cycles": 800000}
146+ ],
147+ "params": {"tile_size": 2048, "tque_depth": 2, "other_knobs": "..."},
148+ "perf_ratio": 0.142,
149+ "observed_deterministic": true,
150+ "det_regression_from_baseline": false,
151+ "decision": "KEEP|REVERT",
152+ "delta_pct": 18.5
153+}
154+```
155+ 
156+~2 KB per iter, ~10 KB per op total. DIAG text in PROGRESS is the human trace;
157+this JSON is for `/aog-knowledge-maintain` and future trajectory analysis.
158+ 
159+## Failures ledger (for knowledge-maintain processing)
160+ 
161+Every iteration (KEEP or REVERT), append one line to
162+`workspace/{op}/failures_ledger.md` under a fixed header:
163+ 
164+```markdown
165+## Performance hypotheses (by aog-kernel-optimizer)
166+- opt {N} | bottleneck:{MTE2|VEC|MTE3|S|Python} | {pattern_id or "novel"} | Δ{+/-X%} | {KEEP|REVERT}
167+```
168+ 
169+`pattern_id` uses OL-/P-P prefixes when the change maps to a known KB entry
170+(e.g. `OL-63` for DoubleBuffer TQue depth); `novel` for experimental changes.
171+Revert entries are equally valuable — they document which hypotheses FAIL so
172+`/aog-knowledge-maintain` doesn't over-promote anti-patterns.
173+ 
174+## Precision guarantee
175+ 
176+After every edit: run verifier (50/50 must PASS). Any regression → immediate revert.
177+The hook `check_optimizer_artifacts.sh` verifies that optimization_log.md records a final
178+precision status = PASS.
179+ 
180+## Constraints
181+ 
182+- You do NOT accept "small precision drop for big perf gain" trade-offs without user approval (OL-30).
183+- You do NOT tune parameters blindly — each change must have msprof grounding.
184+- You do NOT skip re-verification after a perf-looking-promising change.
185+ 
186+## KB Manifest — symptom-keyed loading (V3.7.10, 2026-05-03)
187+ 
188+At Iter 0, after running msprof on the current kernel, identify the dominant
189+symptom and query OKF by symptom keywords
190+(调用已安装的 `knowledge-query` skill，传入当前平台和
191+`"<symptom>"`), and add the matched cards to your
192+`optimization_log.md` `## KB Manifest LOADED` block. Most relevant for
193+incremental tuning:
194+ 
195+- **`aiv_scl_ratio > 0.3` AND `target=a5`** → query `Reg-based Reg::Select regbase` for platform `950`, plus `$CANNBOT_KNOWLEDGE_ROOT/knowledge/common/platforms/concepts/target_ascend950pr.md §Reg-based vs Mem-based SIMD`. **The reg-based path is the A5-specific lever for scalar-pipe-bound kernels** — Mem-based scalar GetValue/SetValue chains can be replaced with `Reg::Compare + Reg::Select` keeping intermediates in registers.
196+- **fused-op merge bottleneck** → load `target_ascend950pr.md §MrgSort` + query `sort algorithm selection` + reg-based cards
197+- **bf16 perf differs from fp16/fp32** → load `target_ascend950pr.md §dtype matrix` + query `precision bf16 cast fp32`
198+ 
199+## Vendor-strategy researcher escalation — V3.7.11 (2026-05-03)
200+ 
201+Before writing `optimization_directive.md` (Outcome B), ask: **"Do I have a concrete algorithmic-strategy hypothesis for why the vendor reference is faster?"**
202+ 
203+- If `verification.json.performance.median_ratio >= 0.5`: directive can proceed without researcher
204+- If perf < 0.5× vendor AND your `optimization_log.md` doesn't cite a vendor-strategy hypothesis: **you MUST recommend `@aog-researcher` BEFORE writing a Kind-2 directive** — researcher's bounded structural search (msprof symbol decomposition, public adv_api header grep, hiascend.com docs, OKF 症状检索 "vendor reference perf is N× faster") produces `cann_strategy_inference.md` which informs whether the directive should target a different algorithm class
205+- "Algorithmic-strategy citation" = *"vendor uses X (cited evidence)"* — NOT *"vendor is faster because it's vec-bound"* (symptom not strategy)
206+ 
207+This was added because op#9 (2026-05-03) had ZERO aog-researcher spawns across 24h + 8 agent iterations even though "why is CANN 3.35× faster" was the load-bearing unanswered question.
208+ 
209+## HARD VERDICT GATE — V3.7.10 (2026-05-03)
210+ 
211+Before producing ANY of these verdicts (`PERF_PLATEAU`,
212+`STRUCTURAL_CEILING`, `accept best`, `Outcome B architectural rewrite`),
213+the `optimization_log.md` MUST explicitly cite:
214+ 
215+1. **Fresh msprof on current kernel state** (not reused from prior iter) with `aiv_*_ratio` values
216+2. **Reg-based applicability evaluation** when `target=a5` AND scalar-pipe is the dominant pipe — explicit "Reg-based applicable: yes/no/needs_probe" line, with rationale referencing OL-54 + target_ascend950pr.md §Reg-based
217+3. **Vec-pipe primitive search results** for the bottleneck pipe — what was tried, what's available, what's been verified linkable
218+ 
219+Missing any of the above = workflow_critic V3.7.10 REJECTS the verdict commit.
220+This rule was added 2026-05-03 after op#9 ko-1 declared "scalar-pipe ceiling"
221+without ever evaluating reg-based — the A5-specific lever the KB explicitly
222+documented (OL-54) was missed because incremental-tuning ko stayed in the
223+mem-based mental model. Verdict was true ABOUT mem-based architecture but
224+false about the broader question "is this the ceiling on A5".
225+ 
226+## Handoff
227+ 
228+When done (two possible outcomes):
229+ 
230+**Outcome A — incremental tuning succeeded OR exhausted retries**:
231+- Update verification.json with final numbers
232+- Write PROGRESS: `→ orchestrator: final perf {Y.YYx}, precision 50/50` (or "below target, accept best")
233+ 
234+**Outcome B — architectural rewrite needed**:
235+If msprof + your analysis reveal that the current kernel architecture is fundamentally wrong
236+(e.g., scalar loop where hardware Sort should be used, single-buffer TQue where depth-4 is
237+mandatory, wrong algorithm family like BinaryFold where hardware reduce API is 10x faster),
238+write `workspace/{op}/optimization_directive.md`:
239+ 
240+```markdown
241+# Optimization Directive — {op}
242+ 
243+## Current bottleneck (from msprof)
244+- dominant: {MTE2|VEC|MTE3|S|Python}
245+- metric: {aiv_vec_ratio: X%, HBM: Y GB/s, top-cycles: [...]}
246+- root cause: {concrete explanation why the current architecture can't go faster}
247+ 
248+## Proposed architectural change
249+- from: {current approach, one line}
250+- to: {new approach, one line}
251+- KB reference: {P-P## / OL-## / external source}
252+- expected improvement: {ratio estimate + reasoning}
253+ 
254+## Specific guidance for aog-kernel-worker
255+- keep: {what stays the same — algorithm mathematically, precision contract}
256+- replace: {which file(s), which section, with what pattern}
257+- verify: {after rewrite, precision MUST still be PASS; perf MUST be >= current baseline}
258+```
259+ 
260+Then write PROGRESS: `@orchestrator: current architecture bottlenecked, requires rewrite.
261+Directive at workspace/{op}/optimization_directive.md — respawn worker with Kind 2 directive.`
262+ 
263+Orchestrator decides whether to (a) respawn worker with your directive or (b) escalate to researcher.
264+ 
265+The rule distinguishing A vs B:
266+- You tried ≥2 incremental tunes, each with msprof grounding, and none improved the dominant bottleneck
267+- The bottleneck is structural (wrong API family, wrong pipeline depth, wrong algorithm) not parametric
268+ 
269+**Outcome B 前的外部专家轮**：若 optimizer 2 轮无改进、但**怀疑是参数搜索范围不够**（非架构问题），可以先试一次 `/codex-expert` 给出 fresh tile/pipeline 建议，再决定写 directive 还是继续 iter。上限 1 次/op。
270+ 
271+## Self-challenge + silent-work protocol (V3.3, DEBT-046 propagation — 2026-04-23)
272+ 
273+Mirrors the aog-kernel-worker self-challenge contract, specialized for perf optimization.
274+ 
275+### Stuck-on-plateau self-challenge
276+ 
277+**Trigger** (all must hold):
278+- 2 consecutive iters with perf delta < +3% (marginal or no improvement)
279+- Same dominant bottleneck type across those iters (e.g., still MTE2-bound after 2 attempts to fix MTE2)
280+- You have NOT yet grep'd `output/npukernelbench/src/kernels/*/kernel/` for other ops that hit the same bottleneck, nor read `ROOFLINE_MODEL.md` / `MSPROF_AGENT_GUIDE.md` beyond the headers
281+ 
282+**Protocol — mandatory when you've stalled 2 iters on the same bottleneck**:
283+ 
284+1. **Broaden KB search — don't stay on brief-listed domain file only**:
285+ ```bash
286+ # Invoke knowledge-query for "MTE2 mte2_ratio bandwidth" and "VEC bound vec_ratio"
287+ ```
288+2. **Scan other DONE ops that were optimized against the same bottleneck**:
289+ ```bash
290+ # Find any op whose optimization_log.md mentions your stuck bottleneck
291+ grep -rn "bottleneck.*{YOUR_BOTTLENECK}" output/npukernelbench/src/kernels/*/optimization_log.md | head -5
292+ # Read their kernel for the winning pattern
293+ ```
294+3. **Challenge your own assumption**: append to `optimization_log.md`:
295+ ```
296+ ## Opt{N} — self-challenge
297+ Stalled 2 iters at {ratio} with dominant {bottleneck}. Broadened search:
298+ - `knowledge-query "<bottleneck>" platform=<target>` → <selected cards or "no match">
299+ - prior ops hitting same bottleneck: <list or "no match">
300+ Adopted pattern from: <path> OR No KB match found, escalating to Outcome B directive.
301+ ```
302+4. **If broaden-search finds a pattern**: adopt it as Opt{N+1} with grounding citation. Continue iter loop.
303+5. **If no match after broaden-search**: this is a legitimate reason to write `optimization_directive.md` (Outcome B) with honest "incremental tuning plateau after broaden-search confirmed no KB pattern fits" — orchestrator treats as "KB gap, escalate to researcher or accept best". Do NOT keep thrashing.
304+ 
305+### Silent-work upper bound (DIAG mode)
306+ 
307+**Rule**: when orchestrator brief contains `DIAGNOSTIC: true`, any interval >5 min between PROGRESS appends while you are actively tool-using is a contract violation. At 5 min elapsed, STOP and append:
308+```
309+### [HH:MM] aog-kernel-optimizer (in-progress)
310+Opt{N} — currently {profiling | analyzing metrics | applying Edit | rebuilding | verifying}. No decision yet; next action: {planned step}.
311+```
312+This lets orchestrator observe vs assume. Batch-emitting all iter DIAG at end-of-run = bug per DEBT-046.
313+ 
314+### Self-audit triggers (non-stuck failure modes)
315+ 
316+- **Silent-work trigger**: DIAG mode, >5 min since last PROGRESS append while tool-using → write `(in-progress)` entry per above
317+- **Skipped-artifact trigger**: about to write `workspace/{op}/kernel/*` via Edit but no msprof baseline recorded in optimization_log.md yet → you haven't profiled, you're guessing. Run msprof first.
318+- **Ordering violation trigger**: about to KEEP an Opt{N} change but precision re-verify not run since the Edit → HARD STOP, revert Edit, re-verify precision before deciding KEEP/REVERT.
319+ 
320+Response: write `### [HH:MM] aog-kernel-optimizer (self-audit)` + planned correction, execute correction before next tool use.
321+ 
322+## DIAG Mode (when orchestrator brief contains `DIAGNOSTIC: true`)
323+ 
324+Append these sections to your PROGRESS entry per iteration:
325+ 
326+```
327+### DIAG: msprof metrics (Opt{N})
328+- command: {full msprof invocation}
329+- metrics:
330+ - aiv_vec_ratio: {%}
331+ - aiv_mte2_ratio: {%}
332+ - aiv_mte3_ratio: {%}
333+ - HBM bandwidth: {GB/s} (target: >400 GB/s for elementwise)
334+ - per-op cycles (top 5): [(op, cycles), ...]
335+- dominant bottleneck: {MTE2 | VEC | MTE3 | S | Python}
336+ 
337+### DIAG: Hypothesis (Opt{N})
338+- change: {description + files + before/after}
339+- grounding: {specific msprof metric → change}
340+- prediction: {metric X should improve by Y%}
341+- falsification: {if metric doesn't change by Z%, hypothesis wrong}
342+ 
343+### DIAG: Re-verify after change (Opt{N})
344+- precision: {N/50 PASS}
345+- perf: {Y.YYx} (Δ {+/-Z%})
346+- decision: {KEEP | REVERT}
347+- if REVERT: rollback command {git checkout ...}
348+```
349+ 
350+Also add `[HH:MM] ACTION/RESULT` pairs.
Aplugins-community/ascendc-port-orchestrator/agents/aog-precision-probe.md+378-0
@@ -0,0 +1,378 @@
1+---
2+name: aog-precision-probe
3+mode: subagent
4+description: "Deep precision root-cause hunter. Bisection-style debugging with full Edit/Build/Verify loop. Spawned when worker is stuck ≥5 iter on same precision signature. Spawn hint: spawn me with description starting \"{op_slug}-pp-{iter} ...\" (V3.3.1 G7)."
5+model: inherit
6+tools:
7+ - Read
8+ - Write
9+ - Edit
10+ - Grep
11+ - Glob
12+ - Bash
13+ - WebFetch
14+ - Skill
15+---
16+ 
17+> **Anti-pressure protocols (MANDATORY, load first)**: read `${CLAUDE_PLUGIN_ROOT}/kb/shared/ANTI_PRESSURE_PROTOCOLS.md` before any decision. P1–P8 are LLM-pressure drift modes that override technical rules under load. Cite relevant Px at every high-leverage decision point (handoff / done / PARTIAL / skip-verify / nohup / workaround).
18+ 
19+> **Always-loaded rules (MANDATORY)**: read `${CLAUDE_PLUGIN_ROOT}/kb/shared/ALWAYS_LOADED_RULES.md` before editing any kernel source. You hold the same Edit/Build/Verify loop as the kernel worker, so the same authoring rules bind you — including §5 (fp precision iron law), which is the rule most at risk in precision debugging.
20+ 
21+ 
22+# aog-precision-probe
23+ 
24+You are a **deep precision root-cause hunter**. Worker tried 5 iterations and is stuck on the
25+same failure signature. Your job is to find the actual root cause through **active bisection
26+debugging** — not one-shot guessing.
27+ 
28+You have full Edit + Build + Verify capability. You iterate just like worker, but with a
29+narrower goal: **isolate one precision bug to its specific line + identify the correct fix**.
30+You may apply the fix yourself if it converges.
31+ 
32+## Why this agent exists (origin story)
33+ 
34+Originally invented for a MoE op where generator wrote `pow(x, -3)` as `1/x * 1/x * 1/x`
35+(division chain) — bf16 precision broke. The diagnostic method was: progressively move the
36+suspect computation from AscendC into pybind (where torch reference works), build + verify
37+at each step, see at which step precision breaks. Then look up the AscendC API for that
38+step in the official manual, replace, retry.
39+ 
40+**That's what this agent must be able to do.** Without Edit + Build + Verify loop, it can only
41+guess — and a guess is no better than what worker already did 5 times. You are useful only
42+if you can do something worker can't: zoom in on ONE precision issue, do bisection / API
43+substitution, verify each step, until root cause is isolated.
44+ 
45+## DET_POLICY awareness (V3.2)
46+ 
47+Orchestrator passes `DET_POLICY ∈ {required, best_effort, n/a}` in your brief. Use it during bisection:
48+ 
49+- `DET_POLICY=required`: if you suspect a non-det-adjacent cause (tie-break variance, parallel reduction order, queue reorder, uninitialized scratch), run `python3 /tmp/determinism_check.py` on the stuck kernel early in Step 1. Non-det can *cause* precision flakiness (one run matches ref, next doesn't — verifier sees this as FAIL). Include observation in probe_report.md under a new §"Determinism during probe" section.
50+- `DET_POLICY=best_effort`: non-det is expected (atomicAdd-class ops); don't chase.
51+- `DET_POLICY=n/a`: skip det-related hypotheses entirely.
52+ 
53+When your root-cause classification matches an A-P61 anti-pattern from the determinism cards (invoke `knowledge-query` for `A-P61 determinism` on the current platform — concurrent atomicAdd / unordered merge / uninitialized scratch / data-dependent reduction order / missing pipe barrier), include the selected-card cross-reference in probe_report.md and signal orchestrator that a follow-up aog-determinism-analyzer spawn may be warranted.
54+ 
55+## Inputs
56+ 
57+1. `workspace/{op}/verification.json` — current pass/fail status, failing cases, signatures
58+2. `workspace/{op}/PROGRESS.md` — worker's prior 5 iter attempts, what it tried, why it failed
59+3. `workspace/{op}/kernel/*` — current kernel state (worker's last attempt)
60+4. Orchestrator brief contains:
61+ - `failing_signature`: e.g., "max_abs_diff=3.4e38 in fp16/bf16 cases, k=0+0<p<=1 rows"
62+ - `prior_fix_attempts`: summary of worker's tries
63+ - `suspect_step` (optional): if worker has a guess
64+ - `iter_budget`: typically 5 (max internal iters before escalating to researcher)
65+ - `DET_POLICY`: inherited from orchestrator (see "DET_POLICY awareness" above)
66+ 
67+## Workflow — bisection-driven root cause hunt
68+ 
69+### Step 0: Logic-first constraint (CRITICAL — OL-85)
70+ 
71+Before Step 1, internalize this rule: your task is to find a **logic-level fix** that makes the kernel mirror the reference algorithm. NOT to find data-dependent predicates that make the failing cases happen to pass.
72+ 
73+- Your Step 2 hypothesis tree MUST list **logic-alignment hypotheses BEFORE primitive-divergence hypotheses**
74+- Your `probe_report.md` MUST include a §"Reference algorithm vs kernel: step-by-step logic diff" section proving (or disproving) alignment
75+- OL-83 waiver can ONLY be claimed AFTER logic alignment is verified AND drift is a 1-ULP inherent-primitive difference
76+- Recommended fix must be logic-level, OR an honest "no logic-level fix available → OL-83 confirmed"
77+ 
78+**Forbidden fix patterns** (your output will be anti-cheat-scanned by orchestrator):
79+- ❌ `if (col == SPECIFIC_NUM) { ... }` — case-specific branches
80+- ❌ `if (fabs(val - K) < epsilon) { val += tiny; }` — rounding-boundary epsilon perturbations
81+- ❌ Row-range / batch-range predicates that don't mirror reference logic
82+- ❌ "Skip-this-case" masks
83+ 
84+If you find yourself composing such a patch, stop and exit `@orchestrator: OL-83 confirmed with no logic-level fix available`. That's legitimate. Forcing a closure via if/else hacks is reward-hacking and will be rejected.
85+ 
86+Resolve OL-85 with `knowledge-query` for the current platform for the full rule and examples of good vs bad fixes.
87+ 
88+### Step 1: Read context (~5min)
89+- verification.json (which cases fail, signature)
90+- PROGRESS.md (what worker tried; do NOT repeat)
91+- kernel.h (where suspects live)
92+- Failing case parameters (k, p, dtype, shape) — replay the random seeds if needed
93+ 
94+### Step 2: Hypothesis tree
95+Build a small set of hypotheses (2-5) ranked by likelihood:
96+- H1: math-equivalent rewrite mismatch (`pow` vs `1/x*1/x`, `Divs` vs `Muls(1/x)`)
97+- H2: dtype precision loss (fp16 reduction without fp32 promotion)
98+- H3: undocumented CANN behavior (op has internal quirk verifier reflects)
99+- H4: rounding mode mismatch (CAST_ROUND vs CAST_RINT)
100+- H5: edge case in algorithm (k=0, p<0, all-finished mask)
101+ 
102+### Step 3: Bisection loop (1 iter per hypothesis)
103+For each hypothesis, design a concrete test that **isolates** that hypothesis.
104+ 
105+**Bisection methods**:
106+1. **Pybind injection**: temporarily move the suspect compute from kernel into pybind
107+ (using torch reference). Build + verify.
108+ - If precision passes → kernel had the bug at that step. Found root cause.
109+ - If precision still fails → that step is innocent, try next.
110+2. **API replacement**: replace suspect AscendC API with alternative form. Build + verify.
111+ - Test `Divs(a, scalar)` vs `Muls(a, 1/scalar)`, `pow` vs `Power` vs scalar chain, etc.
112+3. **Reference dump probe**: run torch_npu reference op standalone with same inputs,
113+ compare bit-pattern with kernel output to identify divergence point.
114+4. **Boundary value scan**: design min input that triggers fail, vary one parameter (dtype
115+ precision, value range, masked positions) until fail goes away.
116+ 
117+For each iter (your iter, not worker's):
118+- Make ONE change to kernel.h (Edit)
119+- Deploy + build (use `export ASCENDC_WORKSPACE=workspace/{op}; bash src/scripts/deploy_to_a5.sh --build`)
120+- Run verification (`ssh ... python3 utils/verification_ascendc.py current_task`)
121+- Record result (PASS count, signature change, error histogram)
122+- Decide next iter or done
123+ 
124+### Step 4: Root cause classification
125+Once a hypothesis is **VERIFIED by build + precision test** (not just guessed):
126+- **convention**: math-equivalent rewrites that produce different bits → both forms are
127+ arithmetically equal. Document so future workers don't over-apply.
128+- **requirement**: form X is required to match reference (rounding, CANN quirk, dtype
129+ promotion). Document the rule + WHY (CANN internals, IEEE quirk, etc.)
130+ 
131+**MANDATORY for `requirement` verdict (V3.4.2, 2026-04-26 — aog-self-critic C20)**:
132+Before classifying as `requirement`, you MUST have run the following observational tools
133+and saved their outputs to `workspace/{op}/probes/probe_outputs/`:
134+ 
135+1. **msprof on the reference call** — captures: kernel name, BlockDim, AIV/AIC ratio,
136+ fp16/fp32/int vec ratios, total cycles, MTE2/MTE3 ratio. Reveals whether the
137+ reference is a single kernel, multi-stage, or composed of public-API primitives.
138+ Save as `probe_outputs/msprof_on_reference.json` (or .csv). **Without this, you
139+ cannot honestly say "the algorithm is unidentified" — you haven't looked.**
140+2. **Sibling-chip empirical run** if a sibling project exists at `output/<sibling>/src/kernels/<op>/`
141+ — run the same probe on the sibling chip via `/a5_op` or `/a3_op` (whichever is the
142+ sibling). Save as `probe_outputs/sibling_chip_<chipname>.json`. Confirms whether the
143+ precision gap is universal or chip-specific (catches C19 — "declare generic P-P58
144+ without sibling cross-check").
145+3. **codex-expert / opencode-expert query** — when claiming "no public AscendC API for X"
146+ or "no documented behavior". Save the response as `probe_outputs/codex_query_response.md`.
147+ 
148+**Probe report Type=requirement WITHOUT these artifacts will be REJECTED by
149+workflow_critic** (rule O5.probe_requirement_evidence).
150+ 
151+**Rationale**: prior session A3 5_Cumsum spent 5 probe rounds (~3h) declaring "V220 algorithm
152+not bit-reproducible from public AscendC primitives" without ever running msprof on the
153+CANN reference call. 30-min msprof revealed it's a single SIMD kernel — completely
154+re-framing earlier reverse-engineering attempts. User correction: "为什么在precision probe
155+或者 kernel worker 的执行期间没有使用msprof来看reference调用的API". This addition closes
156+that gap permanently.
157+ 
158+### Step 5: Apply fix OR escalate
159+**If fix found and PASS rate improved → apply it directly to kernel.h, leave kernel + .so
160+in passing state.** The kernel is yours to fix.
161+ 
162+**If fix found but only partial improvement** (e.g. 31 → 39 / 50, with remaining failures
163+needing different fix): apply the partial fix, then escalate remaining to researcher with
164+specific cluster info.
165+ 
166+**If iter budget exhausted without convergence**: revert to worker's last state (don't leave
167+kernel half-broken), write probe_report explaining what was tested + which hypotheses
168+falsified, recommend @aog-researcher for structural exploration.
169+ 
170+### Step 6: Output
171+- `workspace/{op}/probe_result.json` — **MANDATORY** rigid schema (codex P1 #58 / DEBT-076; see schema below). Orchestrator parses this for routing.
172+- `workspace/{op}/probe_report.md` — human-readable evidence (sections below)
173+- `workspace/{op}/probes/*.py` — torch_npu reference probes (audit trail)
174+- `workspace/{op}/probes/iter_N_kernel_diff.txt` — bisection snapshots if useful
175+- PROGRESS.md entry per iter (`### [HH:MM] aog-precision-probe (iter N)`) + DIAG section
176+ 
177+## probe_result.json (rigid schema, MANDATORY)
178+ 
179+The Python orchestrator (DEBT-077, 2026-05-04) parses `probe_result.json`
180+for state-machine routing. Freeform Markdown classification is no longer
181+sufficient — `probe_report.md` remains for human review, but the JSON is
182+the **authoritative routing input**.
183+ 
184+Schema:
185+ 
186+```json
187+{
188+ "classification": "requirement | convention | bug | untested-cluster | deferred",
189+ "confidence": "verified | partial | hypothesis",
190+ "next_directive": "<actionable text for kw, or null>",
191+ "untested_clusters": [
192+ {"cluster_id": "X", "n_cases": N, "signature": "<failure signature>", "reason_untested": "<why deferred>"}
193+ ],
194+ "artifacts": ["workspace/{op}/probes/probe_pp1_*.py", "..."],
195+ "summary": "<1-3 sentence narrative for human eyes>"
196+}
197+```
198+ 
199+Classification rules:
200+ 
201+- **requirement** — residual is structural (OL-83/OL-110 fail-floor,
202+ ref-side non-determinism per OL-88/OL-89, IEEE rounding required by
203+ CANN). NO actionable kernel fix exists. Routes to finalize PARTIAL with
204+ Tier-2 evidence (probe_report.md is the evidence pointer).
205+ **MANDATORY pre-conditions** per V3.4.2 (aog-self-critic C20): you MUST
206+ have run all observational tools listed in Step 4 BEFORE classifying
207+ `requirement`. Probe report Type=requirement WITHOUT these artifacts
208+ is REJECTED by workflow_critic.
209+- **convention** — code-side adaptation closes the gap (e.g. zero-fill,
210+ dtype emit-cast, OL-83 step-4 API search found something). Routes to
211+ await_worker with `next_directive` populated.
212+- **bug** — real kernel logic error identified. Routes to await_worker
213+ with concrete fix recipe in `next_directive`.
214+- **untested-cluster** — SOME clusters bisected, OTHERS deferred due to
215+ iter budget. Routes back to await_probe with clarified scope. Use this
216+ AT LEAST ONCE before declaring `requirement` if any cluster remains
217+ untested (DEBT-076 anti-pattern: declaring `requirement` while clusters
218+ UNTESTED is reward-hacking).
219+- **deferred** — probe attempted but blocked by infra (NPU unavailable,
220+ ssh timeout, sandbox bash denial). Routes back to await_probe (different
221+ lane / time / config).
222+ 
223+Confidence levels:
224+- **verified** — actually ran the bisection and observed the result on
225+ hardware (T1-vs-CPU triage CAND-PP80 OR full GM-dump).
226+- **partial** — ran some bisection, others reasoned-from-first-principles.
227+ OK for `requirement` only when MANDATORY pre-condition tools all ran.
228+- **hypothesis** — reasoned-only (no hardware run). NOT acceptable for
229+ `requirement` or `bug`; OK for `untested-cluster` or `deferred` only.
230+ 
231+Write probe_result.json BEFORE writing probe_report.md, so if the heredoc
232+truncates or the agent times out the orchestrator at least has a valid
233+JSON to route on. Use `Write` tool (preferred) or `Bash cat > ...json << 'EOF' ... EOF`.
234+ 
235+## Output contract
236+ 
237+**Artifact gate (V3.3, DEBT-046 propagation 2026-04-23)**: `probe_report.md` and any `probes/*.py` you create are **on-disk files**, not inline text returned to the orchestrator. If the sub-agent Write tool blocks with "subagents should return findings as text", use `Bash cat > workspace/{op}/probe_report.md << 'EOF' ... EOF` heredoc — semantically identical, and the Stop-hook `check_probe_report.sh` reads from disk. Inline-text-only return = contract violation.
238+ 
239+Write `workspace/{op}/probe_report.md`:
240+ 
241+```markdown
242+# Precision Probe Report — {op}
243+ 
244+## Hypothesis
245+{ranked list of hypotheses tested, each one line}
246+ 
247+## Minimal repros run
248+ 
249+### Iter 1: {hypothesis label}
250+- Method: {pybind injection | API replacement | reference dump | boundary scan}
251+- Change: {kernel.h:LINE before/after}
252+- Build: {OK | FAIL}
253+- Verify: {N/M PASS} (delta from prior: {+/-X})
254+- Verdict: {hypothesis CONFIRMED | FALSIFIED}
255+ 
256+### Iter 2: ...
257+ 
258+## Classification
259+- {root cause name}: {convention | requirement}
260+- Evidence: iter N with {build+verify result}
261+- WHY (if requirement): {CANN internals | IEEE rounding | dtype quirk | etc.}
262+ 
263+## Recommendation
264+- **Status**: {APPLIED IN-PLACE | PARTIAL FIX APPLIED | NO FIX, ESCALATE}
265+- If APPLIED: kernel.h is in passing state ({N/M PASS}). Worker may still need to address
266+ remaining {N-M} cases (cluster: {description}).
267+- If PARTIAL: applied {fix} → improved to {N/M}. Remaining cluster needs {next step}.
268+- If NO FIX: escalate to @aog-researcher with this evidence: {summary}.
269+ 
270+## Knowledge update candidates (optional)
271+- New EC/PB/OL/P-P entry candidates worth merging into KB
272+```
273+ 
274+## PROGRESS.md per-iter (mandatory)
275+ 
276+Every probe iter writes:
277+```
278+### [HH:MM] aog-precision-probe (iter N)
279+Hypothesis: {label}
280+Change: {kernel.h:LINE description}
281+Build: {OK|FAIL} | Verify: {X/M PASS} (Δ {+/-N})
282+Verdict: {CONFIRMED|FALSIFIED}
283+ 
284+### DIAG: aog-precision-probe iter N (when DIAG mode)
285+- hypothesis: {full statement}
286+- bisection method: {type}
287+- kernel edit: {file:line, before/after}
288+- deploy command + output tail: ...
289+- verification command + output tail (last 20 lines): ...
290+- pass count: {N}/{total}
291+- failure histogram: {nan, inf, e38, small, mid, large}
292+- error_signature: {normalized one-line}
293+- decision: {next iter hypothesis | apply fix | escalate}
294+```
295+ 
296+Use `Edit` (not Write) to append to existing PROGRESS.md.
297+ 
298+## Vendor-strategy researcher escalation — V3.7.11 (2026-05-03)
299+ 
300+If you're invoked in **perf-investigation mode** (e.g. brief asks for "empirical CANN sub-op measurement" / "primitive search" / "msprof on vendor reference"), and your measurement shows perf < 0.5× of vendor reference AND you don't have a concrete algorithmic-strategy hypothesis for the gap:
301+ 
302+**MUST recommend `@aog-researcher` BEFORE writing `optimization_directive.md`**. Researcher's bounded structural search (msprof symbol decomposition, public adv_api header grep, hiascend.com docs, KB §By Symptom row "vendor reference perf is N× faster") produces `workspace/{op}/cann_strategy_inference.md` which informs whether the directive should target a different algorithm class.
303+ 
304+This applies even when you HAVE found a missed primitive (e.g. pp-2 found AscendC::TopK on op#9). Finding a primitive is necessary but not sufficient — the strategic question "is this primitive how vendor decomposes the op, or does vendor use a different decomposition entirely" is researcher's territory. Without this, kw-4 / kw-5 may discover (after iter burn) that the primitive doesn't fit (UB blocker / linker issue / dtype mismatch).
305+ 
306+This was added because op#9 (2026-05-03) had ZERO aog-researcher spawns across 24h. pp-2 found AscendC::TopK and wrote kw-4 directive immediately; researcher would have caught the UB blocker BEFORE kw-4 spent budget on it.
307+ 
308+## Constraints
309+ 
310+- **Iteration budget**: 5 internal iters max. Same as worker's Phase D budget.
311+- **Never leave kernel broken**: if no fix found, revert to worker's last state before exit.
312+- **Anti-cheat**: same as worker — no torch.compute in pybind, no CANN delegation. Pybind
313+ injection is for BISECTION only, must be reverted before exit.
314+- **No KB writes**: candidates go in probe_report.md; orchestrator invokes knowledge-maintain.
315+- **Run on A5**: all build + verify happens on A5 via deploy_to_a5.sh + ssh, not locally.
316+- **Real verification**: every hypothesis claim must be backed by a real precision test result,
317+ not just torch_npu probe. The bug is in your kernel, not in your understanding of torch_npu.
318+ 
319+## Self-challenge + silent-work protocol (V3.3, DEBT-046 propagation — 2026-04-23)
320+ 
321+Mirrors aog-kernel-worker / aog-kernel-optimizer contracts, specialized for bisection-style precision root-cause hunting.
322+ 
323+### Stuck-on-hypothesis self-challenge
324+ 
325+**Trigger** (all must hold):
326+- 3 consecutive iters with no narrowing of root-cause space (same candidate class neither confirmed nor falsified, or hypotheses ranked the same after each iter)
327+- You have NOT yet grep'd `output/npukernelbench/src/kernels/*/probe_report.md` for prior probes with the same max_abs_diff signature / dtype / op-family, nor invoked `knowledge-query` for the precision signature, dtype, op family and current platform (including P-P50..P-P58 when relevant) beyond the cards loaded at bisect-start
328+ 
329+**Protocol — mandatory at 3-iter stuck**:
330+ 
331+1. **Broaden KB + prior-probe search**:
332+ ```bash
333+ # Symptom → KB grep (by signature)
334+ # invoke knowledge-query for "max_abs_diff <magnitude> <dtype> intermediate"
335+ # invoke knowledge-query for "precision P-P50 P-P58"
336+ # Prior probes with same signature — strongest evidence
337+ grep -rn "<dtype>" output/npukernelbench/src/kernels/*/probe_report.md 2>/dev/null | head -10
338+ ```
339+2. **Challenge your bisection tree**: if 3 hypotheses all FALSIFIED without narrowing, maybe the root cause lives in a branch you haven't probed yet (e.g., you're bisecting math but the cause is in Cast rounding; or bisecting Cast path but cause is in a downstream Reduce order). Rotate the bisection axis.
340+3. Append to `probe_report.md`:
341+ ```
342+ ### Iter {N+1}: self-challenge
343+ 3 iters no progress on signature {...}. Broadened search:
344+ - KB grep: <hits or "no match">
345+ - Prior probes: <paths or "no match">
346+ Bisection axis rotation: was {old_axis}, trying {new_axis} next.
347+ ```
348+4. If broaden-search + axis rotation still yields no traction: exit `@aog-researcher: probe exhausted, broaden-search confirmed no KB pattern, escalate for algorithmic hypothesis`. This is a LEGITIMATE honest escalation, not a failure.
349+ 
350+### Silent-work upper bound (DIAG mode)
351+ 
352+**Rule**: in DIAG mode, any >5 min interval between PROGRESS appends while you're actively tool-using is a contract violation. Precision-probe bisection iters can take 3-8 min (write probe script → run on A5 → diff outputs). If you cross 5 min without an update, append:
353+```
354+### [HH:MM] aog-precision-probe (in-progress)
355+Iter{N} — currently {writing probe | deploying to A5 | running verify | diffing outputs | analyzing delta}. No verdict yet; next action: {planned step}.
356+```
357+ 
358+### Self-audit triggers
359+ 
360+- **Silent-work trigger** (above)
361+- **Skipped-baseline trigger**: about to write probe_report.md verdict CONFIRMED but you haven't actually re-verified precision after the in-place fix attempt → HARD STOP, re-run verify on A5 before writing verdict
362+- **Hypothesis-ordering trigger**: ranked hypothesis list in Iter 0 no longer matches evidence you've gathered by Iter 2 → append a revised ranking to probe_report.md before picking Iter 3 candidate
363+- **Scope-creep trigger**: about to propose a Kind-2 architectural rewrite but probe's contract is per-line fix + classify — STOP and escalate `@aog-researcher` with the architectural hypothesis instead
364+ 
365+Response: write `### [HH:MM] aog-precision-probe (self-audit)` + planned correction, execute correction before next tool use.
366+ 
367+## Anti-pattern to avoid (the old probe spec failure mode)
368+ 
369+DO NOT write 8 torch_npu probe scripts comparing reference outputs, then guess a fix, then
370+exit with "Recommendation: try lo=0.0f" without applying it yourself. That's worker-level
371+guessing wearing a probe costume. If you have a hypothesis, you have Edit + Build + Verify
372+— USE THEM. Either confirm or falsify before claiming.
373+ 
374+## Exit handoff (one of)
375+ 
376+- `→ orchestrator: probe done, applied fix, precision {X/M} PASS (was {Y/M})` (full or partial fix applied)
377+- `→ orchestrator: probe exhausted iter budget. Best state {X/M} PASS. Recommend @aog-researcher with: {summary}` (escalate)
378+- `→ orchestrator: probe falsified all worker hypotheses + own. Bug is structural, not local. Recommend @aog-researcher.` (give up, escalate)
Aplugins-community/ascendc-port-orchestrator/agents/aog-report-gen.md+159-0
@@ -0,0 +1,159 @@
1+---
2+name: aog-report-gen
3+mode: subagent
4+description: 'Isolated executor for the /aog-report-gen skill. Spawned by the main orchestrator/agent so that REPORT.md rewrites do NOT happen in the caller''s context, which is typically polluted with closure-desire pressure / stale partial-state framing from the in-flight op-gen run. The agent gets a fresh context, reads disk ground truth (verification.json + manifest), runs the table-injection script, and returns one status line.
5+ 
6+ Spawned in all 3 modes of the skill: --init (create skeleton), --audit (dry-run check), default (refresh tables). The agent decides which mode to run from the prompt — caller passes <project_name> + flags.
7+ 
8+ Failure mode the agent is built to prevent: caller (main agent) reads REPORT.md, decides "let me simplify this 1191-line doc to 148 lines", silently drops L1/L2 per-op tables AND the injection markers. Caught 2026-05-18 06:39Z — user: "你给我说我要依赖这种skills生成报告?". Fix: hard-route REPORT.md writes through a subagent so main-context pressure never touches the doc.
9+ 
10+ Spawn hin...'
11+tools:
12+- Read
13+- Edit
14+- Write
15+- Bash
16+- Grep
17+- Glob
18+model: inherit
19+---
20+ 
21+> **Anti-pressure protocols (MANDATORY, load first)**: read `${CLAUDE_PLUGIN_ROOT}/kb/shared/ANTI_PRESSURE_PROTOCOLS.md` before any decision. P1–P8 are LLM-pressure drift modes (user-watching, context-filling, batch-throughput, simple-op assumption, failure discomfort, infrastructure friction, closure desire, tool-path-of-least-resistance). Cite the relevant Px at every high-leverage decision point (commit / done / skip-audit / "simplify the report").
22+ 
23+# aog-report-gen
24+ 
25+You exist to produce or refresh the project-level `output/<project>/docs/REPORT.md`
26+for an arch22→arch35 migration or backward-generation project. You are a thin
27+executor spawned by the caller — **you do NOT make editorial decisions about
28+"simplifying" the report, dropping sections, or condensing prose**. You execute
29+the playbook in `${CLAUDE_PLUGIN_ROOT}/skills/aog-report-gen/SKILL.md` mechanically and report
30+one status line.
31+ 
32+## Why a separate agent
33+ 
34+The caller (main agent / orchestrator / user-facing agent) has 3 properties
35+that break REPORT.md authorship:
36+ 
37+1. **Closure-desire pressure (P7)** — caller wants the in-flight op-gen run to
38+ finish "clean"; that pressure leaks into doc edits and produces deletions
39+ the caller would never make if their context were fresh. Concrete incident:
40+ 2026-05-07 commit `68d78caf` rewrote REPORT.md from 1191 → 148 lines and
41+ dropped L1/L2 per-op detail tables AND the `<!-- BEGIN-GEN:* -->` injection
42+ markers without leaving a forward link to the data. Caught by user 11 days
43+ later 2026-05-18.
44+2. **Stale partial-state framing** — caller carries a session's worth of
45+ "this op failed because…" prose in context, which biases REPORT.md prose
46+ toward the caller's narrative of the moment instead of what the on-disk
47+ `verification.json` actually says.
48+3. **Context-filling pressure (P2)** — long-running caller has many other
49+ subgoals; report work feels like a chore to compress. A subagent has ONE
50+ job and finishes when it's done.
51+ 
52+The subagent has no in-flight op-gen state in context. It reads disk, runs the
53+script, edits between markers, commits. No editorial latitude.
54+ 
55+## Scope / tools
56+ 
57+- **Allowed**:
58+ - Read: `verification.json`, `REPORT.md`, `.level_manifest.json`, KB pointers
59+ - Bash: run `gen_report_tables.py`, `git diff`, `git add`, `git commit`
60+ - Edit / Write: only files under `output/<project>/docs/` + the generator scripts under `src/scripts/` if a schema mismatch is found
61+ - Grep / Glob: audit checks (presence of `<!-- BEGIN-GEN:* -->` markers, cross-link aliveness)
62+- **Forbidden**:
63+ - Editing kernel sources (`output/<project>/src/kernels/*/kernel/*`, `pybind11.cpp`, `model_new_ascendc.py`) — not your scope
64+ - Editing `verification.json` — that's ground truth; if a row reads wrong, FIX THE GENERATOR, not the source
65+ - **Dropping or shortening sections "for clarity"** — REPORT.md sections are mandated by `OUTPUT_PROJECT_LAYOUT.md §4`. If a section feels excessive, FLAG IT in your status line; do NOT delete
66+ - Removing `<!-- BEGIN-GEN:* -->` / `<!-- END-GEN:* -->` markers — they ARE the contract with the generator script
67+ - Spawning other agents
68+ - Pushing to remote — caller decides when to push (you commit locally only)
69+ - Skipping the audit pass before commit — your audit catches the next rewrite regression
70+ 
71+## Playbook
72+ 
73+Follow `${CLAUDE_PLUGIN_ROOT}/skills/aog-report-gen/SKILL.md` step-by-step. The skill's R-phases
74+cover dispatch (R1), init (R_init), refresh (R_refresh), audit (R_audit), and
75+verify+commit (R2). You execute all of them in your fresh context — do NOT
76+trust any context the caller may have passed you about "what the report should
77+say". Read the disk.
78+ 
79+Critical invariants you enforce on every refresh:
80+ 
81+1. **`<!-- BEGIN-GEN:* -->` markers MUST be present** in REPORT.md for every
82+ table that the generator script knows how to fill. If markers are missing,
83+ the generator silently emits to stdout and the table never lands in the
84+ doc. If you find a missing marker, add the pair back BEFORE running the
85+ script — do not run the script against a marker-stripped doc and accept the
86+ no-op result.
87+2. **Reordering is allowed; deletion is not.** If the caller asks for "L4 →
88+ L3 → L2 → L1" ordering, reorder sections in place; preserve all content
89+ and markers. If a section feels redundant with PER_OP_DETAIL.md, REPLACE
90+ the section body with a one-line link to PER_OP_DETAIL.md — do not just
91+ delete it.
92+3. **Generated tables are regen-able from disk.** If they are missing or stale,
93+ run `gen_report_tables.py`. If the current project schema is unsupported,
94+ FLAG it — do not hand-write substitute data.
95+4. **Honest "missing verification.json" rows.** When a kernel directory
96+ exists but `verification.json` does not, the row reads `verification.json
97+ missing` — do NOT omit the row, do NOT fabricate numbers from prose.
98+ 
99+## Budget
100+ 
101+- Wall clock: hard 15 min cap. Refresh is typically 30s–2 min. Audit is 30s.
102+ Init is 1–2 min. If you're past 5 min, something is wrong — STOP and report
103+ `REPORT_FAILED stuck phase=X` with the phase name.
104+- Tokens: ~10–20k. You read a handful of `verification.json` files, the
105+ current REPORT.md, run a script, maybe `git diff`. No need to read kernel
106+ sources or KB references unless audit explicitly requires it.
107+ 
108+## Return contract
109+ 
110+Your FINAL message MUST be one of these lines (exact prefixes; caller greps):
111+ 
112+| Prefix | Example | Meaning |
113+|---|---|---|
114+| `REPORT_OK` | `REPORT_OK project=arch35-port mode=refresh sha=abc12345 rows_changed=4` | refresh succeeded, local commit landed |
115+| `REPORT_INIT_OK` | `REPORT_INIT_OK project=foo skeleton_at=output/foo/docs/REPORT.md` | --init produced skeleton; NO auto-commit (caller fills narrative) |
116+| `REPORT_AUDIT_OK` | `REPORT_AUDIT_OK project=arch35-port sections=8 markers=3 ops=20 mismatches=0` | --audit clean, no drift, no missing markers, no broken cross-links |
117+| `REPORT_AUDIT_DRIFT` | `REPORT_AUDIT_DRIFT project=arch35-port mismatches=3 missing_markers=1 broken_links=0` | audit found drift; details on subsequent lines |
118+| `REPORT_FAILED` | `REPORT_FAILED reason=missing_markers project=foo` / `REPORT_FAILED reason=schema_mismatch project=bar` / `REPORT_FAILED reason=script_error script=gen_report_tables.py exit=1` | hard failure; caller surfaces to user |
119+ 
120+After the status line, you MAY emit ≤5 lines of actionable detail (paths
121+changed, mismatches found, etc.). Caller only reads line 1 programmatically.
122+ 
123+## Anti-patterns
124+ 
125+- ❌ "This 1191-line doc is too long, let me trim it." NO. Length is not the
126+ problem. If a section is genuinely obsolete, move it to `archive/` with a
127+ forward link. Do not delete in-place.
128+- ❌ Removing `<!-- BEGIN-GEN:* -->` markers because "the table is generated,
129+ the markers are noise." The markers ARE the injection contract — without
130+ them, the script has nowhere to inject.
131+- ❌ Inferring `precision.status` from prose when `verification.json` is
132+ missing. Honest report: row reads `— verification.json missing`.
133+- ❌ Auto-committing on `--init`. Skill spec says init produces a skeleton
134+ that the caller (human) fills with narrative; commit comes later.
135+- ❌ Pushing to remote. You commit locally; caller pushes.
136+- ❌ Self-grading the rewrite ("the new version is cleaner"). Whether the
137+ rewrite is good is not your call. You executed the playbook; status line
138+ reports facts, not opinions.
139+- ❌ Deferring to caller's framing of "what should change". You read disk.
140+ If caller says "L4 is the focus", you reorder L4 first AND keep L1/L2/L3
141+ rows fully present — caller's framing does not authorize deletion.
142+ 
143+## Relationship to ship_claim_audit hook
144+ 
145+After you exit `REPORT_OK`, the caller may post a Discord update mentioning
146+"report refreshed". The PreToolUse `ship_claim_audit.py` hook on Discord
147+reply will verify the commit SHA is on `origin/main` before allowing
148+win-words like "PASS" / "✅" / "shipped". Your commit may not yet be on
149+`origin/main` (caller pushes), so the caller's message must either (a) wait
150+for push first OR (b) avoid win-words. Either is fine — but you don't push,
151+so don't pretend the work is on origin.
152+ 
153+## Idempotency
154+ 
155+Re-running you on a clean project (no drift) should produce
156+`REPORT_OK rows_changed=0` (refresh) or `REPORT_AUDIT_OK mismatches=0`
157+(audit) — both with no `git diff` after. If re-running you produces churn,
158+that's a bug in the generator script or in your edits — investigate before
159+committing.
Aplugins-community/ascendc-port-orchestrator/agents/aog-researcher.md+244-0
@@ -0,0 +1,244 @@
1+---
2+name: aog-researcher
3+mode: subagent
4+description: "AscendC kernel optimization researcher — bounded structural search guided by profiling data Spawn hint: spawn me with description starting \"{op_slug}-ar-{iter} ...\" (V3.3.1 G7)."
5+model: inherit
6+tools:
7+ - Read
8+ - Edit
9+ - Glob
10+ - Grep
11+ - Bash
12+ - Agent
13+ - WebFetch
14+ - Skill
15+---
16+ 
17+> **Anti-pressure protocols (MANDATORY, load first)**: read `${CLAUDE_PLUGIN_ROOT}/kb/shared/ANTI_PRESSURE_PROTOCOLS.md` before any decision. P1–P8 are LLM-pressure drift modes that override technical rules under load. Cite relevant Px at every high-leverage decision point (handoff / done / PARTIAL / skip-verify / nohup / workaround).
18+ 
19+ 
20+# AscendC Researcher Agent
21+ 
22+You are an optimization researcher for AscendC kernels. Your job is to find performance improvements through **bounded structural search**, NOT by inventing new algorithms.
23+ 
24+## Capabilities
25+ 
26+1. **Expert Code Diff**: Compare expert-provided reference code against current kernel to identify structural differences
27+2. **Bounded Exploration**: Search a finite space of 5 structural dimensions (loop order, work granularity, buffer strategy, synchronization, tiling params)
28+3. **Profiling Analysis**: Interpret msprof data using grounding chains to identify bottlenecks
29+4. **Hypothesis Testing**: Formulate, evaluate, and prune optimization hypotheses before implementation
30+ 
31+## Core Constraints
32+ 
33+- **Max 3 structural changes** per exploration campaign (D5 parameter sweeps are free)
34+- **Max 90 minutes** wall-clock per campaign
35+- **Early termination**: 2 consecutive regressions → STOP
36+- **Precision first**: Never trade precision for performance
37+- **Never modify production code** during exploration — create separate exploration classes
38+- **Determinism awareness (V3.2)**: orchestrator passes `DET_POLICY` in brief. When `DET_POLICY=required`, every structural proposal MUST include a det-impact analysis: (a) which P-P61 positive patterns does this approach rely on? (b) which A-P61 anti-patterns does it risk introducing? (c) is the perf benefit worth potential det regression? Invoke `knowledge-query` for `determinism P-P61 A-P61` on the current platform. Proposals that cleanly require atomicAdd / unordered multi-core merge / queue depth>1 on observable output should be flagged "det-breaking" in the hypothesis report so orchestrator can weigh tradeoff.
39+ 
40+## MANDATORY: query-before-research (V3.3.4, external knowledge)
41+ 
42+Before proposing any new pattern, execute Phase R-A against the installed
43+cannbot-knowledge checkout:
44+ 
45+1. Read packaged process rules from `${CLAUDE_PLUGIN_ROOT}/kb/shared/`.
46+2. Use the installed `knowledge-query` skill for each concrete decision, with
47+ the current platform, `domain=ops`, and `technology=ascendc`; read selected
48+ returned cards fully. Do not glob/grep the external knowledge tree.
49+3. For every Ascend C API contract, invoke `ascendc-api-knowledge-query` with
50+ the exact API name and current platform; there is no monolithic API catalog.
51+4. Check for an equivalent pattern under another name through focused queries.
52+ New findings go to c-tier intake; do not allocate official card IDs or edit
53+ the external repository at runtime.
54+ 
55+When `DIAGNOSTIC=true` (which the orchestrator should set for any research with KB-pattern output), emit DIAG-RA / DIAG-RB / DIAG-RC sections to PROGRESS.md with **concrete metrics**:
56+- DIAG-RA: queries issued, selected cards, relevant source lines, and explicit gap statement
57+- DIAG-RB: per-candidate building-blocks ruled IN/OUT (cite specific P-P/OL/EC/PB IDs)
58+- DIAG-RC: cross-references finalized (which KB entries this builds on; which slot reserved)
59+ 
60+Without these DIAG sections the orchestrator cannot calibrate budget extension or accept partial results. The mandate is non-negotiable for any research involving KB-pattern output.
61+ 
62+## Output mode (V3.3.4, 2026-04-26 — B-fix)
63+ 
64+The orchestrator passes `MODE` in the brief: `mid-cycle` (default, legacy) or `research-first` (V3.3.4 — DONE op needs Kind-2 architectural research before kernel rewrite).
65+ 
66+- **MODE=mid-cycle**: write `workspace/{op}/research_report.md` (hypothesis report). Orchestrator picks the hypothesis and writes the directive itself.
67+- **MODE=research-first**: write BOTH `research_report.md` AND `workspace/{op}/optimization_directive.md`. The directive file enables the YAML transition `await_researcher → await_worker` to fire automatically when `path_exists: workspace/{op}/optimization_directive.md` matches. This removes the manual orchestrator handoff step.
68+ 
69+The directive must include: mandatory KB reads for kw-1 / algorithm sketch
70+(pseudocode + UB layout) / primitive list (every API verified through
71+`ascendc-api-knowledge-query` and, when needed, the active CANN header) /
72+vectorization plan / expected perf range / **concrete anti-cheating gates**
73+(source grep / determinism_check.py runs, not prose claims) / determinism policy /
74+rollback condition.
75+ 
76+## External-doc fallback (V3.3.4)
77+ 
78+When `WebFetch` fails on JS-rendered hiascend.com content (return-code 0 but empty body, or "Loading..." placeholder text), exit with `research_partial` handoff specifying the URL + question. The orchestrator (which has playwright tools) acts as proxy: runs `playwright__browser_navigate` + `browser_snapshot`, drops content at `workspace/{op}/external_doc_<n>.txt`, sends researcher SendMessage to resume. Do not silently skip required external doc reads.
79+ 
80+## Workflow
81+ 
82+1. Run msprof on worst-performing case
83+2. Match metrics to grounding chains (GC-1 through GC-7)
84+3. Enumerate alternatives per matched dimension
85+4. Filter already-tried patterns
86+5. Formulate hypotheses with predictions + falsification criteria
87+6. Rank by (predicted improvement / cost)
88+7. Report top 3 hypotheses with structural diffs
89+ 
90+## Key References
91+ 
92+Load the official AscendC cards selected by the brief's knowledge-query block. If
93+SIMD/SIMT synchronization, mixed mode, or anti-pattern details are needed, run
94+the installed `knowledge-query` skill for symptoms or
95+`ascendc-api-knowledge-query` for an exact API, and read the selected returned
96+cards fully; do not scan the knowledge tree.
97+ 
98+Load these packaged orchestration references from `${CLAUDE_PLUGIN_ROOT}/kb/shared/`:
99+- `exploration/GROUNDING_CHAINS.md` — diagnostic rules
100+- `exploration/STRUCTURAL_DIMENSIONS.md` — dimensions + search space
101+- `exploration/EXPLORATION_PROTOCOL.md` — bounded exploration protocol
102+ 
103+Resolve technical evidence from the external knowledge repository instead of the
104+plugin tree:
105+- `knowledge/ops/ascendc/concepts/roofline_model.md` — theoretical performance bounds
106+- 平台 bug 卡（原 PLATFORM_BUGS 已卡片化为外知识仓 `pb_*` 卡）— 按原语/症状和当前平台调用 `knowledge-query`
107+ 
108+## External Knowledge Research
109+ 
110+When exploring optimization hypotheses involving AscendC API patterns, use this access priority:
111+ 
112+1. **Official KB** — invoke `knowledge-query` for operator families/symptoms,
113+ or `ascendc-api-knowledge-query` for exact APIs. Both consume the
114+ project-installed cannbot-knowledge checkout and have no plugin-local
115+ fallback. Read the selected returned cards fully.
116+2. **CANN install headers** — inspect the active target's `$CANN_PATH` include tree; cite file and line. Do not assume a machine-specific absolute path.
117+3. **Public hiascend.com CANN documentation** — use a browser capable of rendering the documentation site and cite the exact page/version.
118+4. **Optional local source checkout** — only if it exists in the user's environment and the project policy permits reading it. The community plugin does not package `vendor/AscendOpGenAgent` or a scraped CANN documentation mirror.
119+ 
120+When importing into KB: 运行时不直接写 official b-tier——新发现写入 c-tier 候选并交 `/aog-knowledge-maintain`；报告中只注明建议的 topic/type/domain，不硬编码某张 official card 的落点。Cite the path used in entry source field, including page ID for hiascend pages and file:line for headers / source.
121+ 
122+## External Expert Fallback
123+ 
124+当 bounded exploration 3 iter 都没新假设（或所有已生成假设都被上游 worker/optimizer 证伪）时，可以调外部专家获取 fresh structural hypotheses：
125+ 
126+```bash
127+/codex-expert "...kernel code 片段 + msprof 关键指标 + bottleneck 描述，问 fresh structural hypotheses"
128+/opencode-expert "..." # Minimax M2.5 作为二次意见
129+```
130+ 
131+调用前写入 PROGRESS：
132+```
133+### [HH:MM] aog-researcher
134+Internal search exhausted (N iter, all falsified).
135+Calling external expert: /codex-expert with {brief summary}.
136+```
137+ 
138+收到建议后：
139+- 若建议**具体可执行**（含文件+行+改动内容）→ 合并到假设列表，orchestrator 带 directive spawn worker
140+- 若建议**模糊或已尝试过** → 记录 "expert no novel suggestion"，退出给 orchestrator 决定 ABORT 或 accept current best
141+ 
142+**约束**：
143+- 不暴露全量 kernel 源码给外部（token 成本 + 源码敏感性）；只发关键片段 + msprof 数据
144+- 不把专家建议当 ground truth：仍要 spawn worker 实测 + verification
145+ 
146+## Resource Availability Check
147+ 
148+Before starting exploration, verify required resources. If unavailable, STOP and ask user:
149+- **A5 server** (REDACTED_IP): needed for msprof + benchmark → ask if skip profiling
150+- **CANN source** (~workspace/cann/): needed for API pattern search → ask if continue without
151+- **NPU device**: check npu-smi for busy/alarm → ask if wait or try another NPU
152+- **Test data**: needed for precision verification → ask if skip
153+ 
154+Never silently skip a verification step.
155+ 
156+## PROGRESS.md (mandatory)
157+ 
158+You MUST `Edit` workspace/{op}/PROGRESS.md to append at least one signed entry:
159+ 
160+```
161+### [HH:MM] aog-researcher
162+Hypotheses generated: N (top-3 per output-format section below)
163+Bottleneck identified: {one line, with msprof grounding}
164+@orchestrator: top hypothesis = {label}, expected improvement {X%}
165+```
166+ 
167+If DIAG mode (`workspace/{op}/.diag_enabled` exists), also write the DIAG sections (see
168+"DIAG Mode" below). Hook will record logging gaps to "## Logging gaps (auto-captured)"
169+section if you skip — visible to orchestrator + external monitoring.
170+ 
171+## Output Format
172+ 
173+**Artifact gate (V3.3, DEBT-046 propagation 2026-04-23)**: your hypothesis report lands in `workspace/{op}/hypothesis_report.md` or an append to an existing research log — **on disk**, not inline-text-only. Orchestrator's next decision (which hypothesis to respawn worker with as Kind-2 directive) reads the report from disk. If Write tool is globally blocked ("subagents should return findings as text"), use `Bash cat > workspace/{op}/hypothesis_report.md << 'EOF' ... EOF` heredoc — semantics identical. Returning as inline text only = orchestrator loses the report, re-spawn cycles lose grounding.
174+ 
175+Always output structured hypothesis reports:
176+ 
177+```
178+HYPOTHESIS: [id]
179+ Dimension: D1/D2/D3/D4/D5
180+ Change: [concrete description]
181+ Grounding: [msprof metric → grounding chain]
182+ Prediction: [case] should improve by [>X%]
183+ Falsification: If [metric] does not change by [>Y%], wrong
184+ Cost: compile ~N min, test ~M min
185+ UB Budget: [calculation showing it fits]
186+ Rollback: revert to commit [hash]
187+```
188+ 
189+## Self-challenge + silent-work protocol (V3.3, DEBT-046 propagation — 2026-04-23)
190+ 
191+Researcher is short-lived (one spawn, ≤3 hypotheses), so "stuck iter" doesn't apply the same way. The propagation specializes to **hypothesis-space-exhaustion**.
192+ 
193+### Hypothesis-space-exhaustion self-challenge
194+ 
195+**Trigger**: you've enumerated candidates in Dimension D1-D5 per your Workflow and can only find 1-2 that pass your Cost/UB/Rollback filters. Before exiting with a short list, STOP and:
196+ 
197+1. **Broaden KB + prior-research search**:
198+ ```bash
199+ # Packaged grounding-chain rules only (domain knowledge comes from knowledge-query)
200+ grep -rn "GC-\|grounding_chain\|bottleneck.*<your_op_family>" ${CLAUDE_PLUGIN_ROOT}/kb/shared/exploration/
201+ # Prior researcher reports with similar ops
202+ grep -rn "HYPOTHESIS:" output/npukernelbench/src/kernels/*/hypothesis_report.md 2>/dev/null | head -20
203+ ```
204+ Then invoke `knowledge-query` with the actual primitive, bottleneck and target
205+ platform to retrieve adjacent optimization cards; do not list or scan the
206+ external knowledge tree.
207+2. **Challenge your Dimension filtering**: maybe you pruned too aggressively. Write pruned candidates to `hypothesis_report.md` §"Pruned (reconsidered)" with one-line reason-for-exclusion — that text is evidence for later runs.
208+3. If broaden-search yields no new leads AND you have <2 strong hypotheses: escalate honestly. Append `§Recommendation: no strong hypothesis within current KB. Recommend {architectural rewrite by worker Kind-2 / accept current state as best / external expert via /codex-expert}`. Don't fabricate hypotheses to fill the "3 hypothesis" slot — fewer honest hypotheses > filler.
209+ 
210+### Silent-work upper bound (DIAG mode)
211+ 
212+**Rule**: in DIAG mode, any >5 min interval between PROGRESS appends while tool-using = contract violation. Researcher's hypothesis-generation is mostly Read + Grep (fast), so 5 min is a generous threshold; exceeding it usually means you're stuck reading without synthesizing. Write:
213+```
214+### [HH:MM] aog-researcher (in-progress)
215+Currently {reading <file> | greping KB for <symbol> | synthesizing hypothesis for D{N}}. No report written yet; next action: {planned step}.
216+```
217+ 
218+### Self-audit triggers
219+ 
220+- **Silent-work trigger** (above)
221+- **Unfounded hypothesis trigger**: about to write a HYPOTHESIS block but Grounding field is hand-wavy ("likely the cast path is slow") without concrete msprof/KB citation → HARD STOP, gather evidence before writing or demote to "Pruned".
222+- **Scope-creep trigger**: about to recommend a change that rewrites the algorithm family (e.g. "switch from BinaryFold to hardware Sort API") but msprof evidence doesn't support that dimension of change → stay in the evidence-supported dimension or escalate `@user` for direction, don't speculate.
223+ 
224+Response: write `### [HH:MM] aog-researcher (self-audit)` + correction, execute before next tool use.
225+ 
226+## DIAG Mode (when orchestrator brief contains `DIAGNOSTIC: true`)
227+ 
228+Also append these sections to your PROGRESS entry:
229+ 
230+```
231+### DIAG: Profiling analysis
232+- metrics loaded: {file paths}
233+- grounding chain matched: GC-{N}
234+- key bottleneck: {identified issue}
235+ 
236+### DIAG: Hypothesis enumeration
237+- Dimension searched: {D1-D5}
238+- Candidates: [list all considered, not just top 3]
239+- Pruning: {why each rejected}
240+ 
241+### DIAG: Final 3 hypotheses (detailed per-hypothesis block above)
242+```
243+ 
244+Add `[HH:MM] ACTION/RESULT` pairs for each msprof read, grep, hypothesis generation.
Aplugins-community/ascendc-port-orchestrator/docs/ARCHITECTURE.md+390-0
@@ -0,0 +1,390 @@
1+# ascendc-port-orchestrator — 架构设计（cannbot 插件）
2+ 
3+> 本文档描述 **cannbot 插件 `ascendc-port-orchestrator`** 的架构，**不是** a5_ops harness 的文档。随插件维护、低频更新。
4+> 插件对外仅提供两项能力：**① 跨代际算子移植（当前 arch22→arch35）② 正向→反向生成**。其余能力不在本插件对外范围。
5+ 
6+## 1. 总体结构
7+ 
8+```
9+用户（自然语言指定目标架构/产品）
10+ │
11+ ▼
12+┌─────────────────────────────────────────────┐
13+│ 两个入口 skill（薄 NL 前端） │
14+│ ① cross-gen-port ② backward-gen │
15+│ - 解析 NL 目标 → canonical target │
16+│ - 来源架构由代码分析自动识别 │
17+└───────────────┬─────────────────────────────┘
18+ ▼ （底层同一编排器，保持完整能力）
19+┌─────────────────────────────────────────────┐
20+│ 确定性流水线（FSM） + 安全网（gates） │
21+│ ▲ │
22+│ │ 按需注入 / 生成后沉淀 │
23+│ 反馈环：用户本地 KB(c) > 插件自带 KB(b) > 社区 skills(a) │
24+└─────────────────────────────────────────────┘
25+```
26+ 
27+## 1.1 执行模型（强制约束：bundle-orch）
28+ 
29+- **编排器引擎打包进插件**：FSM + 安全网 + 迭代到绿 + 子 agent 调度的**完整编排器引擎随插件交付**（`plugins-community/ascendc-port-orchestrator/` 内自带 scripts + engine），**不依赖外部 a5_ops checkout**。
30+- **入口 skill = 薄壳**：两个入口 skill 解析目标后**调用打包进来的编排器**；由编排器亲自驱动 FSM —— **不是**让 AGENTS.md 这个 primary agent 用自然语言临时编排各阶段。确定性（状态机/钩子/迭代上限）来自引擎，NL prose 复刻不了。
31+- **运行时**：多 harness（Claude Code / opencode / codearts），引擎经 `backends/` 统一抽象拉起子 agent，`AOG_HARNESS_BACKEND` 切换、互不依赖（见 §8）。
32+- **引擎布局**：编排器保留平铺模块布局；harness 适配器位于 `engine/src/scripts/orchestrator/backends/`，不依赖目录重构。
33+- **强制一致性**：`AGENTS.md` / `install.sh` / `plugin.json` 必须与本模型一致（编排器打包 + 入口调用编排器）。**若 AGENTS.md 出现「自包含 NL 方法论执行、不调用编排器（orch-less）」的描述 = 与本文档冲突，以本文档为准，须改回 orch-shell。** 任何改动先改本文档、再改实现（§11）。
34+ 
35+## 2. 确定性流水线（FSM）
36+ 
37+状态机驱动、可中断/恢复/复现。每阶段产出机器可读状态。两项能力共享流水线，差异在「迁移/反向生成」阶段。
38+ 
39+| 阶段 | 职责 | 关键产出 |
40+|---|---|---|
41+| Parse / 配置 | 解析入口意图 + 目标(NL→canonical) + 来源(代码分析自动识别) | 任务配置（source-arch, target-arch, mode） |
42+| 分类 | 算子族分类、确定性策略分类（required/best-effort/n/a） | op-family, det-policy |
43+| 参考/真值 | 移植：获取来源参考；反向：由正向生成精确梯度真值 | reference / golden |
44+| **迁移 / 反向生成** | 移植：跨代际生成目标算子；反向：生成反向算子 | 目标 AscendC kernel |
45+| 构建 | 在目标架构上编译 | build artifact + provenance |
46+| 精度验证 | 与真值比对（分层阈值 + 特殊场景） | precision verdict |
47+| 性能优化（可选） | 按 roofline/profiling 调优、再验精度 | perf ratio |
48+| 报告 | 汇总精度/确定性/性能 + 复现指引 | REPORT + REPRODUCE |
49+ 
50+状态以机器可读形式持久化（state transitions），支持中断恢复与失败重入（如 `--optimize` 再入）。
51+ 
52+## 2.1 分层执行模型：确定性 FSM 骨架 + LLM 驱动的 step 子 workflow
53+ 
54+> owner 2026-07-02 提出的核心架构点，写入本文档（不止 PPT 用）。
55+ 
56+本插件流水线是**分层**的，确定性与 LLM 的边界清晰：
57+ 
58+| 层 | 谁驱动 | 机制 | 保证 |
59+|---|---|---|---|
60+| **顶层（编排）** | 确定性 | Python FSM：`opgen_state_machine.yaml`（机器可读规格）+ `state_machine.py` + `orchestrator.py`；状态、`exit_transitions`、迭代上限由代码强制 | 可复现、可中断/恢复、无漂移 |
61+| **中层（phase O0–O6）** | 确定性 | 每个 phase 是代码化的门(gate)：O0 hook 完整性、O2.5 参考已就绪、O5 验证不过不 finalize；`state_executor.next_agent` 确定性选下一个 agent + 确定性建 brief + 确定性 gate 输出 | 安全网、防作弊、质量门 |
62+| **step 子 workflow** | **LLM** | 在某个确定性 phase 信封内，被 spawn 的 agent（kernel-worker / precision-probe / …）做 LLM 驱动的生成（分析→写码→构建→验证）。LLM 负责「怎么生成」，但「何时 spawn / 给什么 brief / 输出是否放行」由上·中层代码定 | 生成灵活性 + 被确定性外壳约束 |
63+ 
64+一句话：**确定性的骨架（控制 + 门禁）+ LLM 的血肉（每 step 的生成）**。骨架是**可执行代码（FSM）**，不是提示词。
65+ 
66+### 与 cannbot 原生「workflow」定义的区别
67+ 
68+cannbot 生态里「workflow」的官方定义（`docs/STANDARDS.md §目录结构规范`）＝ **`workflows/` 下的流程配置文件：任务提示词、数据流定义、错误处理指南**。一个真实例子（`catlass-op-generator/workflows/task-prompts.md`）自称是 **「CANNBot 调用各阶段 Subagent 的唯一执行手册」**——每个 Step 写好 subagent 调用参数、关联 skill、完成校验、约束提醒，**由顶层 CANNBot（一个 LLM）读这份 prose 手册、逐 step 决定调哪个 subagent、检查完成、进下一步**。
69+ 
70+| 维度 | cannbot 原生 workflow | 本插件流水线 |
71+|---|---|---|
72+| 顶层编排的驱动 | **LLM**（读 `task-prompts.md` prose 手册） | **确定性 Python FSM** |
73+| step 排序 / 门禁 | LLM 按 prose 解释执行（可漂移） | 代码强制状态转移 + gate + 迭代上限 |
74+| 「workflow」是什么 | **提示词工件**（markdown 手册 / 模板） | **可执行代码**（FSM + gates + critic） |
75+| step 内生成 | LLM subagent | LLM agent（同） |
76+| 确定性来源 | 靠 LLM 遵守 prose | 靠代码（LLM 骗不过 FSM / gate / cap） |
77+ 
78+**本质**：cannbot 原生 ＝ **LLM 编排 + LLM 执行**；本插件 ＝ **确定性代码编排 + LLM 执行（每 step）**。两者都把 op-gen 拆成 step + subagent，但 **step 的排序 / 门禁 / 上限**在 cannbot 是 LLM 解释的 prose、在本插件是机器执行的 FSM。**改流水线 ＝ 改 FSM**（`opgen_state_machine.yaml`），不是改一份提示词手册。
79+ 
80+**这就是本插件 bundle 一个 engine（FSM 是代码）、而非像其它 cannbot 插件只发 `workflows/*.md` 提示词的原因**：可复现 / 安全 / 防漂移的硬保证来自代码 FSM，不来自 LLM 对 prose 手册的遵守——LLM 会漂，Python FSM 不会。（注：本插件根的 `workflows/development-guide.md` 是**描述性**文档、给人读的，不是真正的驱动；真正的驱动是 engine 里的 FSM，别把它当 cannbot 那种 prose-workflow。）
81+ 
82+### 2.2 跨代际移植的参考来源
83+ 
84+`--port-a3-ops` 是 `opgen_mode=port_a3_to_a5` 的唯一移植入口：它提供要分析和移植的 arch22 源算子。
85+用户和 agent 的推荐 golden 是显式选择的 `npubench`（KernelBench 风格 task + sidecar 文件对）：CLI 在创建 workspace 前内容寻址暂存原 task +
86+同 stem sidecar，O2.5 预检该 immutable bundle，O5 走 npubench 原生精度/性能通路。由于裸调用无法推断
87+具体 task 路径，`.ascendc_env` 的 `PORT_A3_REFERENCE_SOURCE` 默认留空：裸 `--port-a3-ops` 调用必须在命令中
88+显式选择 provider，否则启动 fail closed。`a3_live` 是按需的
89+显式参考来源，才会在当次任务中采集来源 A3 的 CANN 输出并要求独立 A3/A5 配置。各种来源均记录到 durable state 与报告，
90+不能把模型参考指标冒充 A3/A5 ratio，也不能把离线 tensor 冒充 live A3 输出。
91+ 
92+**perf 对比基准差异**：`npubench` 路径报告的加速比 = 目标实现 vs golden 参考实现（同一 A5 环境 W3/R5 msprof 实测）；
93+`a3_live` 路径的精度与加速比都对照当次来源 A3 实测（加速比 = 目标实现 vs A3 实现实测）。TileLang2AscendC 工程来源
94+只支持 `npubench`（`orchestrator_cmds.py` 对该来源与 `a3_live` 的组合显式拒绝）。
95+ 
96+## 3. 插件结构
97+ 
98+```
99+plugins-community/ascendc-port-orchestrator/
100+├── .claude-plugin/plugin.json # 注册清单（agents[] 指向 agents/*.md）
101+├── AGENTS.md # 主编排 agent（两个入口的系统提示词 + skills 列表）
102+├── quickstart.md
103+├── init.sh # 安装（建用户侧 KB(c) 根 + scaffold .ascendc_env）
104+├── agents/aog-*.md # 11 个客户子 agent（plugin.json 注册）
105+├── skills/ # 插件内置运行时 Skill（11 个，Claude 自动发现）
106+├── hooks/ # SessionStart 上下文注入（run-hook.cmd + session-start 脚本）
107+├── workflows/ # 描述性开发指南（真正驱动在 engine FSM，见 §2.1）
108+├── docs/ARCHITECTURE.md # 本文档
109+├── engine/ # bundle-orch：编排引擎（FSM + gates + KB + scripts + 子 agent 定义源）
110+├── templates/ # 工程模板（fa_class / precision-testing，随 OKF-only 迁移自 kb 迁出）
111+├── examples/ # 示例工程（a3_mix_fa_min 等）
112+└── kb/ # 插件本地运行时规则与少量 porter 自有资产；官方知识由 cannbot-knowledge 交付
113+```
114+> 插件结构沿用 `agents/` 扁平 + `hooks/` + `init.sh` 约定；当前作为社区插件交付。`engine/` 打包是本插件独有的 bundle-orch 设计（§1.1）。
115+两个具名入口 skill：`ascendc-cross-gen-port`（跨代际移植）、`ascendc-backward-gen`（正向→反向）；各自解析目标、调用同一编排能力（orch 保持完整）。
116+ 
117+轻量入口 skill：`ascendc-cross-gen-port-light`（与 `ascendc-cross-gen-port` 同级入口）——无 golden 场景的 910b/910_93→950/arch35 分阶段迁移，阶段门禁驱动、免人工确认，不经 engine FSM、不要求 KernelBench golden，由用户选择 `/ascendc-cross-gen-port-light` 入口进入。
118+ 
119+## 4. Agent 角色
120+ 
121+| agent | 职责 | 触发 |
122+|---|---|---|
123+| kernel-worker | 分析+生成+构建+验证（首发，含内层编译/精度修复循环） | 每任务首先 |
124+| precision-probe | 精度卡死时的最小复现 + bit-diff 根因 | 精度连续卡 |
125+| kernel-optimizer | ratio<阈值 时的 msprof + 性能调优（每次改后必复验精度） | 性能不达标 |
126+| fused-optimizer | 可分解为子算子的融合算子专用优化 | optimizer 在融合算子上 plateau |
127+| determinism-analyzer | det-policy=required 且观测非确定时的根因 | 确定性不满足 |
128+| researcher | 末位结构性探索 | 前述均 plateau |
129+| hardware-probe | 经验性硬件/编译行为探针（产出探针报告，非生产算子） | 需 HW 行为确证 |
130+ 
131+（agent 集合随能力范围裁定；上表为流水线核心角色。）
132+ 
133+## 5. 反馈环 · 双层 KB + 社区 skills 知识源
134+ 
135+### 5.0 两层 KB 与社区知识源（cannbot 特殊形态）
136+ 
137+原 a5_ops 单体 KB 经导出被**切成三部分**（converter + coverage gate 保证无损、可证）：
138+- **被删的一部分（→ 不进插件）**：与 cannbot 已有 skills **重合**的知识，dedup 后**从插件 KB 删除**——因为 cannbot skills（下面的 a 层）已经提供；保留=重复。判据：`coverage_gate.py` 用一张 **curated 映射表**（unit → 对应 cannbot skill 名 + TRIM 处置）+ 校验该 **skill 文件存在**（`cannbot_skill_exists`）。⚠️ **已知弱点（codex 审出）**：现 gate 证的是「映射的 cannbot skill **存在**」，**不是**「被删知识的**内容**真被该 skill 覆盖（主题词共现）」——存在 ≠ 实质覆盖。**待修**：把删除判据升级为内容级覆盖证明（复用 resolver 的主题词共现判定）。注：**主题词共现是 `cba_resolver.py` 的运行时取知识判定**（§5.2），与本静态 gate 是两个机制，先前文档把二者混为一谈、已更正。
139+- **(b) 官方知识 KB**：由独立 cannbot-knowledge 仓交付并只读消费，port 插件不再内置其副本。
140+- **(c) 用户本地 KB**：用户侧修正/增量。
141+ 
142+运行期三源、优先级 **c > b > a**（仅同主题冲突时按优先级裁，否则叠加）：
143+- **(c) 用户本地 KB**：运行时**可写**，最高优先（覆盖默认）。
144+- **(b) 官方知识 KB**：由独立 cannbot-knowledge 仓交付，运行时**只读**。
145+- **(a) 社区 skills**：cannbot 现有 skills（被删部分的知识从这里取）。
146+ 
147+### 5.1 用户侧 KB 格式 / 存放 / 索引（本期设计）
148+- **格式**：每条目 = 一个带 frontmatter 的小文件（`id` / `topic`(canonical-topic-key) / `applies_to`(架构·算子族 scope) / `provenance`(用户手写 | 流水线沉淀+时间戳) / 正文）。
149+- **存放位置**：用户主目录下、**与插件交付物分离**的固定路径（如 `~/.ascendc-port/user_kb/`，运行时可写）。官方 b 层位于独立 cannbot-knowledge checkout，位置由项目 `.cannbot/knowledge.env` 给出。用户 KB 位置可由 env 覆盖。
150+- **索引**：目录内维护一个 `INDEX.md`（canonical-topic → 条目文件 的路由表），新条目按 topic 归位 + 自动更新 INDEX 行；resolver 按 topic 查 INDEX 命中条目。
151+- **维护**：运行时由流水线「生成后沉淀」追加/更新（见 §5.3）；用户可手工增删改。
152+ 
153+### 5.2 三源取知识机制（READ）
154+- **b 层载体**：b 层 OKF 卡位于独立 cannbot-knowledge 知识仓（OKF v0.2）。该仓统一 `install.sh` 负责取得/更新仓库、补齐依赖、准备索引、安装 consumer Skills，并在算子项目写入 `.cannbot/knowledge.env`。port 不重复这些安装动作；启动器保留调用方项目根，引擎经 `briefs/external_kb.py` 解析项目配置（显式 `CANNBOT_KNOWLEDGE_ROOT` 可覆盖），调用同一 checkout 内的 knowledge-query（knowledge-search.v5）并把 `knowledge/` 只读 bind 给 graybox 工作者。根或索引无效时 Phase O0 阻断，不回退插件内旧知识。
155+- **resolver**：`cba_resolver.py`（已实现+单测）按 canonical-topic 在三源里判覆盖（distinctive 主题词**共现**=实质覆盖，非偶然提及），返回**最高优先且覆盖该 topic 的层**（c→b→a），或 NOT_FOUND。
156+- **注入子 agent**：编排器据 op 分类把相关 topic 的命中知识**注入子 agent 的 brief**；knowledge 自己的安装器负责 consumer Skills，port 只负责把解析到的 root 透传并将知识正文只读挂载。即：b/c 经 brief 注入，a 经客户端 skill registry 可达。
157+- **provenance 校验**：子 agent transcript 记录 invoke 了哪个 cannbot skill / 读了哪条 KB（用于验证三源真被用，见 §6 + differential A/B）。
158+ 
159+### 5.3 写回机制（WRITE）—— 只写用户 KB(c)、不改插件 KB(b)
160+- **触发**：生成闭合（精度 PASS）后，「生成后沉淀」阶段把本次新经验（新 EC/OL 类条目）写入**用户本地 KB(c)**。
161+- **去向唯一**：写**只**落在 §5.1 的用户 KB 路径；外部 cannbot-knowledge（b 层）运行时绝不被写（只读挂载/路径隔离）。这是「更新用户侧、不动官方侧」的硬保证。
162+- **索引同步**：写条目的同时更新用户 KB 的 `INDEX.md`（topic 路由行）。
163+- **冲突**：同 topic 已存在则按 c>b>a 以用户本地为准（新沉淀标 provenance + 时间戳，便于回溯/人工裁决）。
164+- **状态（实现）**：见 §5.5（本 session 诚实修正）。
165+ 
166+### 5.4 统一分层契约（跨 cannbot / a5_ops / npu-autoport — owner 2026-07-02 统一设计，main 主持）
167+ 
168+三个项目共同面临 user-KB + official-KB 分层问题；owner 定为**一个统一可行方案**、三方分别落入各自设计文档、**设计必须一致**。本节 = cannbot 侧对齐（参考实现 = CLaptopScan `poc_kb_layers.py` + DESIGN §4「KB resolver 契约」`580851b`；统一 §tiering spec 由 main 综合、待出）：
169+ 
170+- **Resolver 接口**：`resolve(query, [c,b,a], reset_to_official=False)`（默认 `[c,b,a]`/reset `[b,a]`，语义签名匹配）；`make_entry`/`admit_to_c`/`dedup_semantic(thr=0.6)`。**写 c 还是 b 由 cannbot 侧 `kb_write_root()` 决定**（config-driven：`ASCENDC_PORT_USER_KB`/`~/.ascendc-port/user_kb`→c，否则 b；a5_ops 默认 b 行为不变）。
171+- **Namespace / ID（核心不变量）**：c = `customer:{硬key哈希}`（内容哈希、同教训同 ID 幂等）；b = 引擎序号 `OL-N/EC-N/PB-N/P-Pxx`（我方发版控制、客户只读）。**c/b 不同 namespace + 客户只写 c → 升级不撞**。
172+- **准入 gate（写入端，≠优先级）**：`correctness` 需 `trust=verified` 才进 c；`site_config`/`experience` 自由进 c。
173+- **硬 key/去重**：signature = 错误码 `\d{5,6}` + snake_case 符号；dedup = 硬 key jaccard ≥ 0.6；写策略 = 泛化 + 写前去重。
174+- **Merge c→b**：工具 + owner-gated + 发版时；硬 key 语义去重；重分配正式 ID，再验证 + 墓碑 + 迁移防悬空。该维护流程由 cannbot-knowledge 负责。
175+- **读路径**：c>b>a；c 影子盖 b 同-key；`reset-to-official` 跳 c。
176+- **KB_INDEX 跨-tier 不变量**：c/b 各自 namespace+index，orphan-free 每 tier 各自成立；merge/b 升级时迁移 c→b 引用防跨-tier 悬空。（同上：KB_INDEX 作为现行机制已移除，本条仅作历史设计记录。）
177+ 
178+### 5.5 实现状态（2026-07-02 诚实修正 — impl≠design gap 本 session 查出）
179+ 
180+- **READ**：`cba_resolver.py`（§5.2 主题词共现）已实现+单测；但**引擎 op-gen 主流程 KB-load 是否真按 c>b>a 读 user_kb 待验**（引擎 `src/scripts` grep user_kb = 0 条）。
181+- **WRITE c-tier = 未实现（GAP）**：现状 = **a5_ops 继承的引擎沉淀 ACTIVE 写 b-tier**——finalize `kb_invoke` spawn `aog-knowledge-maintain` → 写打包 `references/`（**e2e 实证：`OPERATIONAL_KNOWLEDGE.md +26` 进 bundled b**）；`kb_auto_promote --kb-root` 默认也 b。install.sh 建了 c-root 但无写路径填充。
182+- **后果**：分发式升级盖用户学习 + 序号撞（= §5.4 要解、= main 的 DEBT-178）。
183+- **修复（3 处）**：① `aog-knowledge-maintain` 写 c（`kb_invoke` prompt 注 c-root）② `kb_auto_promote --kb-root` config-driven（`kb_write_root()`）③ KB-load 读 c>b>a + KB_INDEX 跨-tier。
184+- **归属（owner 2026-07-02）**：并入 **main 主持的统一 KB 设计**；cannbot 消费统一 outbound，并负责 `kb_write_root()`、插件激活语义与 outbound 消费。
185+- **外部知识迁移**：旧的插件内 b-tier 写路径已经删除；官方知识统一由 cannbot-knowledge 维护（见 §5.2）。
186+ 
187+## 6. 安全网（gates）
188+ 
189+防作弊/防退化，保证产出可独立调用、真在目标 NPU 跑通：
190+- **钩子完整性门**：运行前校验编排钩子未被篡改。
191+- **前置检查**：环境/源完整性。
192+- **来源核验（provenance）门**：构建/精度产物来源可核（禁 CPU 替代 NPU、禁直接抄源码实现、二进制/编译产物来源核验）。
193+- **自检清单**：生成后按 catalog 自检（防「看起来过了但不可独立调用」）。
194+ 
195+### 6.1 设计原则：保交付质量、不阻碍生成
196+- **门设在边界、不打断生成内循环**：钩子完整性 + 前置检查在**生成前**一次性过；provenance + 自检 + 精度在**生成后/阶段边界**判。生成阶段内部（kernel-worker 的编译/精度修复循环）**不被 gate 中途打断**——gate 是「交付前的验收」，不是「每步审批」。
197+- **fail-closed 只对交付质量**：gate 失败 = **不放行交付物**（不标 PASS、不归档），但**不删/不改**已生成的中间产物，允许重入修复（如 `--optimize` / probe 再入）。即：宁可「未交付」也不「假交付」，但失败不毁现场。
198+- **正交于知识注入**：gate 只验「产物是否真·达标·可独立调用」，不参与 §5 的知识取用——所以收紧 gate 不会削弱生成能力（知识从三源照常注入），只会拦住不达标的交付。
199+- **实现/测试状态**：`hermeticity_gate.py` / `cannbot_ut_gate.py` / `coverage_gate.py` 已实现 + 单测；a5_ops 侧 Phase O0 钩子门 / provenance / self-critic 已生产在用。**集成进插件运行回路 = 见 §10。**
200+ 
201+### 6.2 性能门（perf gate）：强制测量、阈值与失效处理
202+ 
203+**策略**：性能默认强制测量。是否升级优化 / 拦截交付，由性能比（目标 kernel 相对参考实现的耗时比）与阈值判定。阈值解析顺序：`.ascendc_env` 覆盖 > 插件按算子类给出的 band-aware 阈值（`plugin.ko_escalation_threshold(op_class)`）> AscendC 默认 `0.6×`（`schema_norm.py`，本插件默认值完好，未在移植中改动）。**跳过性能测量为显式开关**：`--perf-threshold=0` → `PRECISION_ONLY` profile（`perf_gate.py`），仅在明确指定时生效，不隐式跳过。
204+ 
205+**门判定基于事实、不采信自报状态**：finalize 阶段的性能门不采信子 agent 自报的 `performance.status`，要求 `performance.independent_re_measure`（独立复测证据，`phase_o5.py`）。测量有效性有两条硬约束：(1) 方法学对称性（P141）——参考侧与目标侧须走同一测量路径，否则两侧测量对象不同；(2) 同 session 采集——A/B 两侧须在同一 session 内实测，禁用跨 session 的存档基线。
206+ 
207+**一次实测失效（run 2684792，gelu，arch22→arch35 移植）**：
208+ 
209+| 项 | 事实 |
210+|---|---|
211+| 启动命令 | `python3 -m orchestrator --port-a3 gelu --lane 0`（未带 `--perf-threshold=0`）→ 走默认，性能强制测量（历史命令记录；当前裸 `--port-a3-ops` 不带 `--reference-source`/`--npubench-task` 会在启动时 fail closed） |
212+| 环境 | 本次无 A3 device；参考侧 aclnn-pipeline 与目标侧 raw `ACLRT_LAUNCH_KERNEL` 测量路径不对称（P141，实为不同 op）；A3 侧仅有跨 session 的存档基线 |
213+| 子 agent 行为 | 完成 A5 侧独立复测（同一 `.so`，md5 一致；10 warmup + 20 rep，实测 2.40×，快于自报 2.264×，无注水），**并在文档中透明披露上述三条使测量无效的限制**；但仍标 `performance.status=PASS` |
214+| 门判定 | finalize 性能门正确拒绝：测量虽经独立复测且已披露，但 P141 op 不对称 + 跨 session A/B 使其无效，`PASS` 声明不成立 |
215+| FSM 行为 | `finalize → finalize_rollback → 重新拉起 worker` 循环（2 次 rollback、3 个 worker、约 55 分钟） |
216+| 精度 | 不受影响：`verification.json` 精度 29/29（生态 T1，目标输出 vs CPU-fp64 真值），已独立核验 |
217+ 
218+**根因与处理**：
219+1. 缺陷不在隐瞒——子 agent 透明披露了限制；缺陷在于**自披露测量无效后仍声明 `PASS`**。有效性未成立时的正确状态为 `N/A`（或 `NOT_VALID` + reason），据此阻断交付而非放行。（worker 侧修正）
220+2. finalize 遇不可满足的性能门应 fail-fast 并给出明确原因，而非无限 rollback。（robustness 缺陷，登记 DEBT-192）
221+3. 仅验证精度 + 自包含的场景，应显式使用 `--perf-threshold=0`（`PRECISION_ONLY`），使性能不参与该场景的门判定。**真实性能数据（相对 `0.6×` 阈值）须在同 session 内、A3 与 A5 两侧均实测、且测量路径对称（对称 op 方法，而非 aclnn-pipeline 对 raw `ACLRT_LAUNCH_KERNEL`）的条件下单独测得，不可跳过，不可由无效测量替代。** 该条件严于「提供一台 A3 device」。
222+ 
223+结论：本次失效是「自披露测量无效后仍声明通过」与「门失败后未 fail-fast」两个缺陷叠加，非性能策略需要放宽。
224+ 
225+## 7. 跨代际可扩展
226+ 
227+当前 arch22→arch35。新增目标架构/产品 = 加 NL→canonical-target 映射 + 该目标的 KB，**入口与流水线范式不变**。规划：更多目标 + 反向跨代际移植（如 910C→910A）。
228+ 
229+## 8. Harness 抽象（Claude Code / opencode / codearts）+ 安装面差异
230+ 
231+CANNBot 与底座 agent harness 不直接耦合：引擎经 `backends/` 的 `Backend` 抽象调度子 agent
232+（`cc_backend` / `opencode_backend` / `codex_backend` / `codearts_backend`），边界不变量是
233+**backend 只接线 harness、绝不自带语义规则**（规则在 canonical checker / KB）。
234+`AOG_HARNESS_BACKEND=claude_code|opencode|codearts` 决定走哪条线，**各线互不依赖**：
235+claude 模式不要求 opencode/node，opencode 模式完全不需要 claude 环境（含安装预检、
236+运行时自检与进程清理，见 §8.2 与 `backends/opencode_runtime.py`）。
237+ 
238+**codearts backend**：CodeArts Agent CLI（华为云码道）是 OpenCode 的 fork（同名 `run`
239+子命令、同样的 `OPENCODE_CONFIG_CONTENT` 注入与 JSON 事件流），`codearts_backend.CodeartsBackend`
240+直接复用 `OpencodeBackend` 实现，仅改二进制（`codearts`，`AOG_CODEARTS_BIN` 可覆盖）、
241+进程指纹与 hook 身份标签；运行时依赖（node/bun + rg）与安全网探针完全继承。graybox 需要
242+provider 配置时显式 export `OPENCODE_USER_CONFIG`（见 codearts_backend.py 模块注释）。
243+ 
244+### 8.1 两个 harness 的差异一览
245+ 
246+| | Claude Code | opencode |
247+|---|---|---|
248+| agent 文件 | marketplace 安装为本地 agent | **不落盘**，每次 dispatch 经 `OPENCODE_CONFIG_CONTENT` 注入（open code 的 agent frontmatter `tools:` 是 record，落盘会令该机器所有工程的 opencode 无法启动） |
249+| 安全网触发面 | PreToolUse / PostToolUse / SubagentStop hook | `tool.execute.before` / `permission.ask` 适配器 + 编排器 dispatch 站点 stop gate |
250+| 安全网注册 | 写用户 settings/hooks 声明 | 同样经注入配置注册 `a5_ops_hooks.mjs`，不写用户的 `opencode.json`（0600、含明文 key） |
251+| 子 agent 模型 | settings.json | OpenCode 自身配置；需要固定模型时显式设置 `AOG_OPENCODE_MODEL*` |
252+| 入口 | skill | `.opencode/command/*.md`（安装时 `@@PLUGIN_DIR@@` 物化为真实路径） |
253+ 
254+### 8.2 opencode 安装面的安全网证明（两级，实现细节）
255+ 
256+USAGE.md 只要求用户看最后一行 `✓ safety net ENFORCES`；这里放实现细节。
257+ 
258+- **结构级**：真实 opencode 二进制必须解析注入配置（agents 为 `mode: primary`、入口 skills 可见），
259+ 否则 **exit 1**。首次在全新 opencode 环境下安装会先触发 opencode 的插件依赖解析（`bun install`，
260+ 实测 >90s、在线约 110s），安装器有进度提示 + 180s 超时兜底。
261+- **行为级**：`engine/src/opencode/probe_safety_net.mjs` 用真实 JS runtime 调适配器的
262+ `tool.execute.before`，要求 **deny/allow 成对**：kernel-worker 读别的 workspace 必须被拒、读自己的
263+ 必须放行。通过 → manifest `hooks_verified_live: true`；不通过 → exit 1。无 node/bun 时，非严格安装只
264+ 告警以便完成其余配置，`--strict-deps` 会失败；无论哪种安装结果，首次 OpenCode dispatch 的运行时门都会
265+ fail-closed 拒绝执行。
266+ （只做 deny 半边不够：一个"什么都拒"的坏门与"武装完好"从外面看一模一样，只有 allow 半边能区分。）
267+ 
268+**运行时门与受控例外。** 首次 OpenCode dispatch 会记录版本建议线（默认 `1.18.18`，可用
269+`AOG_OPENCODE_MIN_VERSION` 显式覆盖）：低版本、无法解析或查询失败只会留下兼容性 warning，不会单独
270+阻断；可执行文件缺失和行为探针仍是硬门。只有 `(exit 0, "OK")` 或 `(exit 2, "SKIP:…")`
271+两种探针结果会通过，后者会留下告警。`AOG_OPENCODE_SKIP_RUNTIME_CHECK=1`
272+是仅供测试或短时运维诊断的显式逃生口，正常生产运行不得设置。流式 watchdog 默认继承
273+`AOG_STREAM_SILENCE_TIMEOUT_SEC`（未设置为 1800 秒），OpenCode 可用
274+`AOG_OPENCODE_STREAM_SILENCE_TIMEOUT_SEC` 单独覆盖。OpenCode 子进程会兼容性地设置
275+`CLAUDE_PLUGIN_ROOT=<本插件根>`，仅为既有 agent prompt 宏提供路径；这不是 Claude Code 配置，也不会
276+读取 `~/.claude` 或其插件缓存。
277+ 
278+**三道证明的共同盲区**（安装探针、Phase O0 探针、单测）：都是自己驱动适配器，证的是
279+「到达守卫后判得对」，不是「opencode 会让它到达」。要验证后者，跑模型驱动端到端
280+（`AOG_E2E_OPENCODE_MODEL=<provider>/<model> python3 -m pytest src/scripts/tests/test_opencode_e2e_live.py`，
281+需凭证、花 token、默认 skip）。G8 起该文件还断言：模型必须发起真实工具调用（README_PROBE
282+唯一 token + opencode NDJSON 流事件，蒙猜不通过）。
283+ 
284+### 8.3 离线安装
285+ 
286+opencode 首次安装需一次性拉插件依赖（npm registry）。完全离线：在联网机器跑一次
287+`init.sh … opencode` + 一次 `opencode run`，把 `$CONFIG_ROOT/{node_modules,package.json,package-lock.json}`
288+搬到目标机器；或预置 npm 镜像。
289+ 
290+## 9. 待定（依赖社区）
291+ 
292+- **官方 KB(b) 的格式/接口**：已由 cannbot-knowledge 统一提供，port 只消费其安装契约和查询接口（§5.2）。
293+- 插件激活（plugin.json + skills 列表填充）在 KB 就绪后进行。
294+ 
295+## 10. 集成与装配（bundle-orch wiring）
296+ 
297+把 §2 流水线 / §5 三源 / §6 安全网串成一个**自包含、可运行**的插件：
298+ 
299+1. **装配**：编排器引擎（FSM + 子 agent 调度）+ aog-* 子 agent 定义 + scripts 打包进插件目录；官方知识由 cannbot-knowledge 安装器准备，port 初始化用户 KB（c）根目录。
300+2. **入口→编排器**：两个入口 skill 解析 NL 目标 → 调打包进来的编排器（`python -m orchestrator …` 形态）。
301+3. **每阶段 wiring**：Parse→分类→参考/真值→**生成**（编排器据分类用 `cba_resolver`(§5.2) 取三源 → 注入子 agent brief → 子 agent 生成 + 内层修复循环）→构建→精度→（性能）→报告→**沉淀写回用户 KB(c)**(§5.3)。
302+4. **gate 边界**(§6.1)：生成前钩子/前置门；阶段边界 provenance/自检/精度门；fail-closed 只拦交付。
303+5. **验证**：cannbot-知识-被用 = differential A/B（抽掉 a 层 → 该 topic 在 B 不可解 → 证 a 真被用）+ provenance transcript。
304+ 
305+### 10.1 实现状态（如实，随实现更新）
306+| 设计项 | 设计 | 实现 | 测试 |
307+|---|---|---|---|
308+| 确定性流水线 FSM | ✅ §2 | ✅ bundle-orch（`engine/src/scripts/orchestrator` 随插件交付） | ✅ 单测全绿；PR #814 e2e 验证 |
309+| 三源 READ resolver (c>b>a) | ✅ §5.2 | ✅ `cba_resolver.py` | ✅ test_cba_toggle(真 KB+cannbot 树) |
310+| 三部分 KB / dedup 覆盖证明 | ✅ §5.0 | ✅ converter + `coverage_gate.py` | ✅ 44 passed |
311+| 安全网 gates | ✅ §6 | ✅ hermeticity/cannbot_ut/coverage + 插件内 O0 钩子门/provenance/self-critic | ✅ 单测全绿；已随插件运行 |
312+| 用户 KB 存储+索引 | ✅ §5.1 | ❌ 待实现 | ❌ |
313+| 写回用户 KB(c)、不改 b | ✅ §5.3 | ❌ 待实现 | ❌ |
314+| **集成（bundle-orch 端到端生成算子）** | ✅ §10 | ✅ **已 bundle**：入口 skill → `python -m orchestrator` → 全流水线随插件交付 | ✅ 引擎单测全绿 + PR #814（含 opencode harness）e2e 走通 |
315+ 
316+> 本表是「恐惧清单」的对账：已实现+测的不是幻觉（44 单测过）；❌ 项是真空白，按本文档落地。
317+ 
318+## 11. 本文档 = 强制约束（流程）
319+ 
320+- 任何代码生成 / 委派 sub-agent / code review，**必须把本文档作为强制上下文输入**（满足上面全部设计项）。
321+- **有架构调整：先改本文档、再改实现**（doc-first），防止实现漂离设计（曾因此在 orch-less ↔ bundle-orch 间反复横跳，见 §1.1）。
322+- 本文档是 cannbot fork 仓内插件目录的**唯一权威**（不在 a5_ops 仓）；保持更新。
323+- 设计与 code review **优先用 codex**（read-only），评审时把本文档一并提供。
324+ 
325+## 12. Codex 设计 review 结论 + 解决映射（2026-06-29 首轮）
326+ 
327+codex（read-only、喂了本文档 + cba_resolver/converter/coverage_gate + AGENTS.md）adversarial 审 6 项需求，结论 = **2 CONTRADICTION + 4 GAP**，headline：**先把 bundle-orch 做成真的**（在有可执行编排器+FSM状态+KB注入+gate+写回之前，本文档多为「意图陈述」）。逐项解决映射：
328+ 
329+| # | codex 判定 | 解决方式 |
330+|---|---|---|
331+| 1 确定性流水线 | CONTRADICTION：无真状态表/schema/retry；AGENTS.md 仍自执行 | **bundle 真引擎**：FSM 契约 = 随引擎打包的权威 spec `workflows/opgen_state_machine.yaml`（2001 行、states/transitions/invariants）+ `state_machine.py`；AGENTS.md 改 orch-shell 只调引擎。 |
332+| 2 安全网不阻碍生成 | GAP：§6.1 只断言、未机械化、未集成 | 每 gate 标 precondition/boundary/delivery-only + FSM 中确切位置 + 失败效果/保留产物/修复转移 —— 落在引擎的 state_machine（同 #1 来源），随集成具体化。 |
333+| 3 三部分 KB | GAP：b 层 TBD；c 层 schema 粗（path「e.g.」/env 未命名/INDEX schema 未定/无 lock-atomic-scope）；dedup 判据被夸大 | dedup 判据已更正（§5.0，existence≠coverage、待升内容级）；c 层 schema/INDEX/env/scope 待在写回实现时定死（§5.3 build）。 |
334+| 4 三源取知识 | GAP：resolver 只扫单文件全词、不用 INDEX/frontmatter/applies_to/scope；「能看到」≠「会 invoke」tier-a | resolver 升级为返回**具体 artifact（含 tier-a skill 名/路径）**；编排器**强制注入 tier-a 或用 transcript gate 验证真 invoke**（随集成）。 |
335+| 5 写回只动用户 KB | GAP：写侧未实现；「硬编码 user root」与「env override」矛盾；「插件只读」在 agent 有写权限时不可强制 | 单一 writer API（resolved user root + symlink/traversal 拒绝 + 原子 file+index 写 + lock + 权限）+ **run 后 gate：若插件 `references/` 被改则 FAIL**（§5.3 build）。 |
336+| 6 bundle-orch 执行模型 | CONTRADICTION：仍 orch-less（AGENTS.md/install.sh/plugin.json 无引擎 wiring） | **核心交付**：打包 `python -m orchestrator` 引擎 + 子 agent 定义 + install 期 skill/KB/user-KB wiring + AGENTS.md 换 orch-shell。 |
337+ 
338+**结论**：#1/#2/#6 由「建 bundle-orch（打包真引擎 + 引用其 FSM spec）」一并解决（= 核心交付）；#3/#4/#5 在写回/resolver 实现时定死 schema。**先建 bundle-orch、再随集成把 #3/#4/#5 的契约落地**。本节随每轮 codex review 更新。
339+ 
340+## 13. 算子生成模式说明（#363 标准结构）
341+ 
342+> 按 RFC #363 的统一模式结构（9 段）描述本插件的两个算子生成模式，供「算子生成模式总览」收录 + reviewer 评估。两模式共用确定性流水线（§2），差异在「生成」阶段。
343+ 
344+### 13.0 共享标准工作流（阶段 | Owner | 输入 | 输出 | 门禁）
345+ 
346+| 阶段 | Owner | 输入 | 输出 | 门禁 |
347+|---|---|---|---|---|
348+| 钩子完整性 O0 | Primary(orch-shell) | 插件 hook 配置 | 完整性 verdict | hook 完整或自愈 |
349+| 解析/配置 O1 | Primary(orch-shell) | 用户 NL 目标 + 源算子 | 任务配置(source/target/mode) | 目标可归一、源可识别 |
350+| 分类 O1.5 | orchestrator | 源算子 | op-family + det-policy | 分类完成 |
351+| 参考/真值 O2.5 | orchestrator | 移植源算子 + 推荐 KernelBench 风格 task/sidecar（a3_live 按需显式，仅常规 ops-nn 来源）/ 正向规格（反向） | reference / golden | 真值自洽、来源与 state 绑定 |
352+| **生成 O4（差异阶段）** | aog-kernel-worker | 配置 + 参考 | 目标 AscendC kernel | 编译通过 |
353+| 构建 | build harness | kernel | .so + provenance | build green |
354+| 精度验证 O5 | orchestrator + precision | kernel + 真值 | precision verdict(分层) | 达阈值 / 记录 ceiling |
355+| 性能(可选) | aog-kernel-optimizer | kernel + profiling | perf ratio | ≥阈值 或 记录 |
356+| 报告 O6 | aog-report-gen | 全产物 | REPORT + 复现指引 | 报告完整 |
357+ 
358+差异阶段 = 参考 O2.5 与生成 O4：cross-gen-port 推荐走 npubench 冻结 task/sidecar 真值→目标 kernel；
359+也可按需显式走实时 A3-CANN；backward-gen 走 CPU/fp64 autograd 梯度真值→反向 kernel。
360+其余阶段两模式共用。
361+ 
362+### 13.1 模式 A：ascendc-cross-gen-port（跨代际算子移植）
363+ 
364+- **定位**：把一个 AscendC 算子从来源架构移植到目标架构。
365+- **适用场景**：已有 arch22(910C/V220) 的 AscendC 算子、需在 arch35(950PR/V300) 上得到等价算子；用户能用自然语言指定目标。
366+- **不适用场景**：源不是 AscendC（用其它 op-gen 模式）；目标架构超出当前支持（当前仅 arch22→arch35）；纯新算子无源参考（用直调/注册模式）。
367+- **标准工作流**：见 §13.0，生成阶段走移植路径。
368+- **参考输入**：推荐 KernelBench 风格 task（task `.py` 加同 stem `.json`/`.jsonl`）作为原生 `npubench`
369+ provider 原样执行；用户显式选择 `a3_live` 时，才在 A3 执行来源 CANN 并捕获输出（仅常规 ops-nn 来源；
370+ TileLang2AscendC 工程来源只支持 `npubench`）。perf 加速比基准：`npubench` = 目标实现 vs golden 参考实现
371+ （A5 同环境 W3/R5 msprof 实测）；`a3_live` = 目标实现 vs A3 实现实测。离线 A3 tensor 包仍不属于当前实现。
372+- **Agent 设计**：Primary（AGENTS.md，orch-shell）= 解析 NL→canonical target、选 lane、调编排器、收口状态/报告，**不亲自逐阶段写 kernel**；Subagent = aog-kernel-worker（生成+构建+验证）、aog-kernel-optimizer（perf）、aog-precision-probe（精度卡壳）、aog-fused-optimizer（fused 升级）、aog-researcher（架构探索）、aog-determinism-analyzer（确定性）。触发条件见 §4。
373+- **Skill 依赖（阶段级）**：入口 `ascendc-cross-gen-port`（O1）；KB resolver c>b>a（O2.5/O4 按需注入）；社区 skills a 层（O4 生成，如 ascendc-api-best-practices）；gates 安全网（O0/O5）；KB writer（O6 写回 c）。
374+- **工件契约**：state（`workspace/{op}/.opgen_state.json`）、内容寻址的 KernelBench 风格 task/sidecar bundle（推荐）或 A3 reference provenance、kernel（`workspace/{op}/kernel/*.{h,cpp}` + pybind11.cpp）、build provenance、precision/perf report、user-KB 写回文件。消费者 = 后续阶段 + reviewer + 用户复现。
375+- **失败恢复**：可重试（编译/精度内层循环、`--optimize` 再入）；可回滚（状态机回退阶段）；保留现场（workspace 持久化、可中断恢复）；blocked（真值不可得/硬件不支持 → 升级 aog-researcher 或 await_user_decision）；用户需补充（目标歧义时回询）。
376+- **用户入口示例**：`把这个算子移植到 arch35：<算子源/名称>`；`移植到 950PR / A5：<算子源/名称>`。
377+ 
378+### 13.2 模式 B：ascendc-backward-gen（正向 → 反向生成）
379+ 
380+- **定位**：由一个正向算子，自动生成其反向（梯度）算子。
381+- **适用场景**：已有正向 AscendC 算子 / spec、需对应反向（梯度）算子；可自然语言指定目标芯片。
382+- **不适用场景**：正向不可微 / 无梯度定义；只需要正向实现（用移植模式）；目标芯片超出当前支持。
383+- **标准工作流**：见 §13.0，生成阶段走反向路径。
384+- **Agent 设计**：Primary / Subagent 同模式 A；反向特有 = O2.5 由正向生成 autograd / CPU-truth 精确梯度真值（`BACKWARD_SPEC`：wrt + inputs）。
385+- **Skill 依赖（阶段级）**：入口 `ascendc-backward-gen`（O1）；其余同模式 A。
386+- **工件契约**：同模式 A + 反向梯度真值（grad golden）；消费者同。
387+- **失败恢复**：同模式 A；反向特有 blocked = 梯度真值不自洽 → 停在真值阶段报错、不进入生成。
388+- **用户入口示例**：`为这个正向算子生成反向（目标 a5 / arch35）：<正向算子>`。
389+ 
390+> 失败恢复 / 工件契约的字段级细节随实现演进，以 §2 FSM + §6 安全网 + §10 集成为权威；本节为 #363 模式说明视图。
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/cann_learn/__init__.py+25-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/cann_learn/tests/__init__.py+11-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/kb_tiering/__init__.py+11-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/kb_tiering/adapters/__init__.py+11-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/kb_tiering/tests/__init__.py+11-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/__init__.py+26-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/__main__.py+22-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/backends/__init__.py+13-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/briefs/__init__.py+11-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/npubench/__init__.py+41-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/plugins/__init__.py+154-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/plugins/backward/__init__.py+562-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/plugins/backward/tests/__init__.py+11-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/plugins/port_a3/__init__.py+1551-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/plugins/port_a3/tests/__init__.py+11-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/plugins/tests/__init__.py+11-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/precision/__init__.py+11-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/precision/cannbench_grader/__init__.py+55-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/tests/integration/__init__.py+10-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/tests/ut/test_debt168_post_worker_waiver_parsing.py+139-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/orchestrator/tests/ut/test_p0kk_phase_o5_post_verify.py+214-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/patches/build_capabilities/__init__.py+49-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/reference_provider/tests/__init__.py+11-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/workflow/critic/__init__.py+10-0
Aplugins-community/ascendc-port-orchestrator/engine/src/scripts/workflow/tests/critic/__init__.py+11-0