2026-11-09 统一入口与公共核数 helper（52a99e4）

• 对应评审 199300662，expand 和 head mix 的前后向在 5f331a2 已统一 eager 与编译的注册算子入口。薄 wrapper 继续负责原始标量校验和有效核数传递。
• 对应评审 199302053，将 resolve_vector_cores 和 validate_vector_cores 移到 src/cann_ops_tilelang/common/compile_utils.py，由两处算子模块共同导入。采用仓库支持的命名空间目录，避免新增包初始化文件触发全量选测。

移动后导入及 setuptools 包发现检查通过，仓库选择器确认增量选中 6 个相关 MHC 文件。Ascend auto/PTO 各 66 项受影响测试通过，覆盖注册、Fake、原始标量、组合重放和配置切换。Ruff、197 文件格式和差异检查通过。下文各 123 项为统一入口改造时的完整回归。

2026-11-09 统一 Custom Op 调用入口（5f331a2）

按 PR #61 更新公共入口。eager 和 torch.compile 共用注册算子，公共 wrapper 处理注册前的校验、输出契约与实际配置。注册后由 _impl 执行真实计算，Fake 按原 schema 推导元信息。

• expand.py 中，expand 前后向把有效向量核数传入注册算子，保留空输入短路及调用方缓冲区写入。
• head_compute_mix.py 中，前向在进入注册算子前检查 mhc_pre_eps、token_block_size 的原始类型。反向校验公开标量，再传递按输入设备解析的有效核数。前后向均返回 None。
• test_mhc_graph_configuration.py 中，Fake 隔离测试改为拦截当前使用的核数解析入口。

NPU 非 benchmark 回归 auto/PTO 各 123 项通过。每后端 21 份证据、76 条图事件记录，确认实际重放且无 eager 回退。新进程导入、Ruff、197 个 Python 文件格式及 diff 检查通过。

以下小节保留此前提交的验证记录。本轮运行范围与结果见上文。

本轮 CodeCheck 整改（7dcb63d，2026-11-05）

针对run#59报告的21项告警整改，本地验证已完成，线上CodeCheck需本提交重跑后确认；Ruff通过不等于线上扫描通过。

本次PR最终为7个文件：两个算子模块、common目录下的compile_utils.py、三份测试文件和一份测试helper。没有夹带其它批次，未改共享config.py、architecture.md或两个kernel，未调整精度/fullgraph/缓存上限/mutation。

验证：auto/PTO各 123 passed /16 benchmark deselected，CPU 25 passed；fixture及命名空间最终调整后auto再验 53项入图测试通过；注册/Fake预检7项通过；Ruff及 197个Python文件格式、diff检查通过。原有40条断言、公有入口和4个Custom Op签名均经审计保留。历史性能数据未在本次重测，不将静态整改声明为性能提升。

PR功能描述 / 为什么需要这个合入

为 expand_to_mhc_forward、expand_to_mhc_backward、mhc_head_compute_mix_forward 和 mhc_head_compute_mix_backward 添加 torch.compile 入图支持，使它们能够参与 npugraph_ex 固定 shape 推理图的捕获与重放。

四个入口保留调用方传入输出/梯度缓冲区的接口及原有 eager 行为：正确声明实际写入参数，不新增 optional out，不改变返回 None 的约定。适配覆盖只读输入、输出覆盖、参数梯度分片归并、空输入和配置变化；原始标量类型在进入 Custom Op schema 前校验，避免 bool/int/float 转换导致 eager 与 compiled 接受范围不同。

本 PR 仅包含计划 PR-B 的四个入口，不包含 Engram、norm_fn、MHC pre/post/Sinkhorn 或多层重计算适配。

各文件修改

注册契约与 Fake

Fake 只使用 Tensor 元信息与传入的有类型核数，不启动 kernel、不读取 Tensor 值、不查询硬件或检查环境文件。真实实现保留运行环境与存储别名检查。向量核内部参数检查整数类型与正值；head backward 同时校验调用方核数、有效上限和两项参数梯度分片的 shape。空 head backward 保持将参数梯度分片清零的原行为；空 expand 不执行设备任务，不以空分支宣称捕获重放成功。

配置追踪与工具链兼容

• 删除旧批量快照的 _config_key 字符串及私有配置属性依赖，只传实际参与执行或有效性校验的整数。expand 的值贯穿 wrapper→Custom Op→impl→现有 JIT kernel；head backward 的 kernel 仍使用调用方显式核数，内部上限用于同一轮合法性校验。
• 通过公开 get_num_ai_cores 保留设备选择和 override 语义；编译分支在算子内部读取环境键存在性和值，保留有效范围检查。实测当前 torch_npu 的 os.getenv 日志补丁在设置核数环境变量时会触发 fullgraph 断图，故避免在编译追踪中调用该补丁。
• 参照当前 master，只有 require_ascend 的环境可用性检查登记为编译常量；实际实现仍执行该检查。可变核数不使用 assume_constant_result，配置测试在同一 compiled 对象内切换和恢复，不清 JIT 缓存或重建对象强迫生效。
• 为新增/修改文件补完整 MIT 文件头，生产文件保留来源 LICENSE 的 2026 DeepSeek 版权。未改共享 config.py、docs/architecture.md 或两个 kernel 文件。

验证

基于 master 2ded9ae0648ef754a0de8c47f54838b4fba5fa18，最终单提交 f3da407。

这些是本批独立工作区本次运行的结果，不引用之前24算子快照的通过数量。线上CodeCheck/SCA/UT需PR发布后确认。

运行范围为原有 test_expand.py、test_head_compute_mix.py 与新增三份测试文件，非 benchmark 共 123 项（57 项原有、66 项新增），16 项 benchmark 未运行。新增图用例保持 backend='npugraph_ex', fullgraph=True, dynamic=False，检查逐轮真实 acl_graph_capture / acl_graph_replay 及 fx_run_eagerly=0。独立参考沿用原精度：expand 逐项相等，head atol=1e-5, rtol=1e-4；原有测试及精度标准不变。

核数切换在子进程内验证，避免旧配置模块没有状态快照/恢复接口而污染父进程；同一个切换序列内不清 Dynamo/JIT 缓存，切回原值要求复用原图。记录 kernel 实际收到的 num_vec_cores：expand 随 16/20 改变，head backward 保持 caller 的 4，变化的是有效上限。Fake 测试将运行时核数查询替换为抛异常的函数，确认使用显式元信息即可完成验证。

CI 的依赖分析在本地实际运行，结果为 incremental，选中上述 5 个测试文件；没有修改选择器、CI 配置、测试级别、benchmark 或显存画像。

实际测试命令（已加载下述 CANN/配套工具链，分别设置 TILELANG_DEFAULT_TARGET=auto 和 pto）：

OPS_TILELANG_TEST_LEVEL=0 OMP_NUM_THREADS=1 OPENBLAS_NUM_THREADS=1 \
python -m pytest tests/mhc/test_expand.py tests/mhc/test_head_compute_mix.py \
  tests/mhc/test_expand_compile.py tests/mhc/test_head_compute_mix_compile.py \
  tests/mhc/test_mhc_graph_configuration.py -m 'not benchmark' -x -v --tb=short

实机环境与支持边界

• Ascend 950DT_9574，CANN 9.3.0，Python 3.12.3，torch 2.10.0+cpu / torch_npu 2.10.0。
• TileLang 0.1.15+cpu.git194c1b89，配套 TileLang PTO overlay、PTOAS VMI 0.1.9 和 helper compatibility commit 7c7032b55ecc6b9dfa7c4c9412c93d587d4d3547；auto/PTO 均在已加载 CANN 的配套环境运行，不代表任意默认 PTO 安装已验收。
• 固定 shape、显式前后向推理调用；图测试使用一致 torch.inference_mode()，autograd 数值参考在其外运行。没有新增自动 autograd、.backward() 图捕获、动态 shape 或多设备验收声明。
• 本次未重测性能。此前批量快照的单算子测试中这些入口慢于直调，不能用入图功能通过宣称本提交已获得单算子或整网加速；更完整性能收益仍需另行测量。

AI assistance: OpenAI Codex / GPT-6，负责封装、整改、测试和描述；范围及提交由人工指定并授权。