engram_grad_w_reduce 的 TileLang 直调入口尚未注册 Torch Custom Op，无法作为稳定算子节点接入 torch.compile。本 PR 接入 npugraph_ex 固定形状 ACL Graph，并保留 master 的五个参数、原地累加与返回 None 的行为。

接口与实现

前三个张量只读；两个 FP32 梯度缓冲区按原公式连续累加：

grad_weight_hidden += grad_w_partial.sum(0) * weight_embed.float()
grad_weight_embed += grad_w_partial.sum(0) * weight_hidden.float()

注册 cann_ops_tilelang::engram_grad_w_reduce，明确声明 mutates_args=('grad_weight_hidden', 'grad_weight_embed')。公共入口在编译时调用注册 op，普通执行沿用真实实现；没有新增或恢复 out 参数。

各文件修改

未修改 kernel 数学逻辑或调度、共享 config.py、docs/architecture.md、benchmark 基线或显存画像。

正确性与契约验证

编译入口使用 torch.compile(..., backend='npugraph_ex', fullgraph=True, dynamic=False)，调用处统一使用 torch.inference_mode()。

• 入图形状 (num_partials, hidden_size)：(1,256)、(7,256)、(37,2560)、(38,3072)、(75,6144)、(7,4096)，hc_mult=4，覆盖多个原 kernel 调度分支。
• 从非零梯度缓冲区开始，连续验证首次调用、原存储更新全部只读输入、只读输入更换地址；随后单独更换持久梯度缓冲区。
• 每次根据当前输入演进独立 PyTorch reference；两个梯度均满足原标准 calc_diff < 1e-10，未放宽精度。
• 检查前三个输入不变、返回 None、梯度缓冲区地址不变；使用带非零 storage offset 的连续视图，检查前后未使用区域未被修改。
• 组合模型验证上游缩放、原地更新及下游读取/求和；opcheck 检查实际注册 op 的 test_schema / test_faketensor。
• eager 和 Fake 均验证零分片、错误 hidden/hc/rank/dtype/shape/布局的拒绝行为；原接口不接受空分片，未改变该限制。

Ascend NPU 验收结果

实际命令（脚本还设置 CANN 环境、线程数及 PYTHONPATH）：

OPS_TILELANG_TEST_LEVEL=0 python -m pytest tests/engram/test_engram_grad_w_reduce.py -m 'not benchmark' -x -v --tb=short
python -m pytest -m 'not npu and not benchmark' -x --tb=short
python -m ruff check .
python -m ruff format --check src tests
git diff --check

ACL Graph 捕获与重放证据

auto/PTO 分别在独立进程使用非空 (7,4,256) 输入逐次采集 CPU profiler，数值在每次调用后验证。两条路径结果相同：

固定状态地址时稳定重放；更换被修改缓冲区后合法重新捕获，新缓冲区实际写入正确，没有 eager 回退。以上是图执行取证，不是性能计时，本 PR 不声明加速比例。

软件环境与支持边界

Ascend950DT_9574（28 Cube / 56 Vector），CANN 9.3.0，Python 3.12.3，torch 2.10.0+cpu，torch_npu 2.10.0，TileLang 0.1.15+cpu.git194c1b89，容器 tl_wjx。

auto 使用原 TileLang 安装。原容器默认 PTO 工具链存在 ptodsl._allreduce 兼容问题；PTO 验收使用隔离安装的官方 ptoas-vmi 0.1.9，以及官方 TileLang #3469 的三个 Python helper overlay（固定 commit 7c7032b55ecc6b9dfa7c4c9412c93d587d4d3547）。未覆盖容器原安装，也未将第三方工具链修改加入本 PR。PTO 通过结论仅适用于这套配套环境。

本次验收固定形状的显式梯度归约入口；未声明自动 autograd、动态 shape、高阶梯度或端到端训练支持，性能待单独验证。

完整命令、环境导入路径、日志和逐次图计数保存在服务器 /home/w00947321/aclgraph-experiments/engram-grad-w-reduce-20261030。

CodeCheck 修复（2026-10-31）

初次门禁仅报测试访问 Tensor 受保护属性 _base（G.CLS.11）。改为在创建偏移视图时显式保存底层缓冲区，并将其传给边界校验 helper；保留梯度缓冲区前后未使用区域检查，不改算子或 kernel，不放宽精度。修复后 auto / 配套 PTO 各重新通过 21 项非 benchmark 测试、Ruff、192 文件格式检查及 git diff --check。修复合并为原 PR 的单提交 8f030ff。

修复后线上门禁

流水线 #26 已全部通过（2026-10-31）：CodeCheck pass；线上 UT auto 27 passed，PTO 27 passed，ut.sh 退出码 0。已核实检出的合并提交包含源提交 8f030ff8889b9f03e73e12dc687e6aa3bd0477b0。线上 UT 使用 CI 自身工具链，其结果独立于本地隔离 PTO 环境验收。