已合并
fix: 修复门禁静态检查与开源相似度告警。 #24
fix: 修复门禁静态检查与开源相似度告警。 #24
已合并
machenggui-coder创建于 26 天前
共 99 个文件变更+6259-1554
@@ -41,3 +41,11 @@ AGENTS.md
41# 测试报告导出件(不纳入版本管理)41# 测试报告导出件(不纳入版本管理)
42test_result/42test_result/
43_errors/43_errors/
44+extra-info/
45+ 
46+# examples / ST 运行产物
47+examples/**/build/
48+examples/aclnn/**/output/
49+tests/st/**/build/
50+tests/st/**/input/
51+tests/st/**/output/
@@ -6,45 +6,26 @@
6 6 
7### Added7### Added
8 8 
9-- 合入 origin/master 的融合注意力算子:`FusedBiasSoftmax` / `FusedSoftmaxGrad`(含 ophost tiling UT、NPU 精度对比与 Python API `oil_gas_ops_prospect.fused_bias_softmax`)。9+- 融合注意力算子 `FusedBiasSoftmax` / `FusedSoftmaxGrad`,以及组合接口 `attention_core`。
10-- AscendC 的 Python API 统一放在 `ascendc/pybind/<op_family>.py`,`ascendc/pybind/` 打包为 `oil_gas_ops_prospect` 包根。10+- 统一 Python 包 `oil_gas_ops_prospect`:AscendC 经 `ascendc/pybind/` 暴露,Triton 暴露 `real_rfft` / `real_irfft`。
11-- 合入算子开发规格(现位于 `docs/zh/develop/operator-dev-spec.md`)。11+- 根目录三种构建入口:`bash build.sh` 只编译出包,`--pkg` 只出 OPP `.run`,`--install` 安装并冒烟。
12-- 补齐开源合规交付件:`CONTRIBUTING.md`、`SECURITY.md`、`OAT.xml`、`Third_Party_Open_Source_Software_List.yaml`、`Third_Party_Open_Source_Software_Notice`。12+- 统一测试入口 `bash build.sh -u`(切分、上板、精度、性能、示例、FFT、文档门禁);`--opensource-report` 生成开源测试报告。
13-- 新增顶层 CMake 构建体系:`CMakeLists.txt`、`version.cmake`、`cmake/third_party/googletest.cmake`、`cmake/security_flags.cmake`。顶层 `cmake -S . -B build` 可直接配置并运行 ophost 单元测试,C++ 目标统一应用安全编译选项。
14-- 新增 `.gitcode/workflows/` 门禁流水线配置:编译(x86 / aarch64)、无卡 LLT(ophost UT + Triton CPU 回退精度 + Triton AOT 编译检查)、静态检查。
15-- 新增 `docs/` 文档目录,含快速入门、目录结构、算子列表、接口列表、开发与调试指引。
16-- 新增 `examples/` 示例代码与 `scripts/` 辅助脚本(许可证头检查)。
17-- 新增 `requirements.txt`、`CHANGELOG.md`、`.clang-format`、`.pre-commit-config.yaml`。
18-- `build.sh` 新增 `--ut` 模式,一条命令跑完全部无卡单元测试。
19-- 新增 `tests/README.md`(用例分类与 NPU 依赖说明)与 `tests/conftest.py`。
20-- 新增 `ascendc/pybind/`,集中存放各 AscendC 算子的 Python API 与 C++ 绑定;`pybind_extensions.py` 按 `ascendc/pybind/csrc/<op_family>/*.cpp` 自动发现并生成扩展模块 `<op_family>_npu`。
21 13 
22### Changed14### Changed
23 15 
24-- 根目录 `build.sh` 按 [ops-math `scripts/build_options.sh`](https://gitcode.com/cann/ops-math/blob/master/scripts/build_options.sh) 拆出选项解析:新增 `scripts/build_options.sh`,无卡 UT 主入口改为 `-u`(`--ut` 仍可用),支持 `--ophost` / `--triton` / `--aot` 子集、`--noexec`、`--cov`、`--gtest_filter`。16+- AscendC OpType 去掉 `Og` 后缀(`ComplexMul` 等);Python 接口为蛇形命名,不加 `npu_` 前缀。
25-- 仓库许可证由 CANN Open Software License Agreement 2.0 改为 Apache License 2.0(`LICENSE`、源文件头、`OAT.xml`、`pyproject.toml`)。17+- `operators/` 只保留实现;测试在 `tests/`,aclnn 样例在 `examples/aclnn/`。
26-- 根目录测试目录由 `test/` 更名为 `tests/`,构建脚本、CI、CMake 与文档同步更新。18+- 许可证改为 Apache License 2.0。
27 19 
28-- 去掉 AscendC OpType 的 `Og` 后缀:`ComplexMul` / `ComplexMulGrad` / `FusedBiasSoftmax` / `FusedSoftmaxGrad`,aclnn 同步为 `aclnnComplexMul` 等。20+### Fixed
29-- AscendC Python 公开接口为蛇形命名,不加 `npu_` 前缀:`complex_mul`、`fused_bias_softmax`、`fused_softmax_grad`、`attention_core`。21+ 
30-- 根目录构建入口 `build_install_ops.sh` 更名为 `build.sh`,与 CANN 社区仓惯例对齐。原 `ascendc/build.sh` 保持不变,仍是 AscendC 侧的分步入口。22+- `complex_mul` 的 Python autograd 接线,反向梯度可正确回传到输入。
31-- **测试用例统一收敛到根目录 `tests/`,算子目录不再包含测试用例**:`ascendc/tests/ophost/` → `tests/ascendc/ophost/`;`ascendc/tests/complex_mul/` → `tests/ascendc/complex_mul/`;`triton-ascend/real_fft/` 下的 `test_triton_fft_real.py` 与 `verify_triton_fft_real_128_3d.py` → `tests/triton/real_fft/`。23+- `bash build.sh -u` 单 flag 入口无法进入测试的问题。
32-- `ascendc/build.py ut-ophost` 在检测到 CMake 缓存中的源码目录与当前不一致时自动清理重配,避免搬迁后复用旧缓存报错。
33-- pybind 封装从测试目录移出:C++ 绑定在 `ascendc/pybind/csrc/`,构建脚本为 `ascendc/pybind_extensions.py` / `ascendc/pybind_setup.py`。
34-- **删除源码树中的 `oil_gas_ops_prospect/` 目录**。Python 对外 API 按算子名暴露,不再带 `ascendc` / `triton` 二级包:`oil_gas_ops_prospect.complex_mul`、`oil_gas_ops_prospect.real_fft`。`setup.py` 用 `package_dir` 在打包时组装;裸 checkout 由 `scripts/local_pkg.py` 提供同一导入名。
35-- ophost 单元测试的 googletest 依赖从 `github.com` 直连改为 gitcode 三方件镜像,并抽取到 `cmake/third_party/googletest.cmake` 供顶层与独立配置两种方式复用。
36-- `README.md` 重写,补齐概述、版本配套、源码下载、目录结构、贡献指南 / 安全声明 / 许可证 / 所属 SIG 链接与问题反馈入口。
37-- 按现有实现补齐算子文档:`docs/zh/op_list.md` 为每个 AscendC / Triton 算子写明数学定义、输入输出与约束;`docs/zh/api_list.md` 补全 `fused_*` / `attention_core` / `make_basis` 参数表与 aclnn 两段式入口;`ascendc/README.md` 登记融合注意力算子及其切分/测试说明。
38-- 刷新算子开发规格并迁到 `docs/zh/develop/operator-dev-spec.md`:对齐当前 `tests/`、`ascendc/pybind/`、`build.sh` 与命名约定,去掉开源重整前的 `python/`、`testplan/`、`delivery/` 等过时要求。原 `docs/operator-dev-spec.md` 保留跳转。
39 24 
40## [0.1.0] - 2026-08-2725## [0.1.0] - 2026-08-27
41 26 
42### Added27### Added
43 28 
44-- 首次提交 oil-gas-ops-prospect 油气勘探领域昇腾自定义算子库。29+- 首次发布油气勘探昇腾算子库。
45-- AscendC 算子 `ComplexMul`:复数逐点乘,float32 / ND,AIV-only,产出 CANN OPP `.run` 包。30+- AscendC:`ComplexMul` / `ComplexMulGrad`(float32 / ND,产出 OPP `.run`)。
46-- AscendC 算子 `ComplexMulGrad`:`ComplexMul` 的反向算子。31+- Triton-Ascend:`real_rfft` / `real_irfft`(`norm="ortho"`,支持 autograd)。
47-- Triton-Ascend 算子 `real_rfft` / `real_irfft`:实数正向 FFT 与半谱重建,支持 autograd,非 NPU 环境回退到 PyTorch。
48-- 基于 msopgen 的多算子构建流程(`ascendc/build.py` + `ascendc/build.sh`)。
49-- ophost 侧 tiling 单元测试(googletest),无卡可跑。
50-- 统一 Python 包 `oil_gas_ops_prospect` 与 whl 打包配置。
@@ -19,7 +19,7 @@
19#19#
20# AscendC 算子本身(op_host / op_kernel)不由本文件编译:CANN 自定义算子必须由20# AscendC 算子本身(op_host / op_kernel)不由本文件编译:CANN 自定义算子必须由
21# msopgen 依据当前 CANN 与 SOC 生成配套工程后再交给 opbuild / bisheng 编译,入口是21# msopgen 依据当前 CANN 与 SOC 生成配套工程后再交给 opbuild / bisheng 编译,入口是
22-# `bash build.sh` 与 `bash ascendc/build.sh --pkg`,详见 ascendc/README.md。22+# `bash build.sh --pkg` 与 `bash ascendc/build.sh --pkg`,详见 ascendc/README.md。
23 23 
24cmake_minimum_required(VERSION 3.16)24cmake_minimum_required(VERSION 3.16)
25 25 
@@ -50,18 +50,21 @@ ascendc/operators/${op_name}/
50ascendc/operators/manifest.tsv # 追加一行(tab 分隔)注册算子50ascendc/operators/manifest.tsv # 追加一行(tab 分隔)注册算子
51ascendc/pybind/${op_name}.py # Python API(autograd + 回退)51ascendc/pybind/${op_name}.py # Python API(autograd + 回退)
52ascendc/pybind/csrc/${op_name}/bindings.cpp # C++ 绑定(打进 whl,非测试资产)52ascendc/pybind/csrc/${op_name}/bindings.cpp # C++ 绑定(打进 whl,非测试资产)
53-tests/ascendc/ophost/ # 追加 tiling UT,并更新 CMakeLists.txt53+tests/ut/op_host/${op_name}/ # tiling UT(自动汇入 ophost_ut)
54-tests/ascendc/${op_name}/ # NPU 精度/性能对比测试54+tests/st/aclnn<OpType>/ # 多 case 两段式 aclnn + 真值
55+examples/aclnn/${op_name}/ # 两段式 aclnn 小样例
56+tests/ascendc/${op_name}/ # NPU Python 精度/性能对比
55```57```
56 58 
57-测试用例统一放在仓库根目录的 `tests/` 下,算子目录只放实现与交付件。59+C++ 测试与样例放在仓库根 `tests/`、`examples/`;`operators/` 只放实现。
58 60 
59约定:61约定:
60 62 
61- **切分算法必须抽到 `op_host/<op>_tiling_core.h`**。`<Op>TilingFunc` / `InferShape` 的签名依赖 `gert::TilingContext`,CANN 二进制包未提供 gert faker,无法直接 UT;把与平台无关的算法抽出来后,ophost UT 才能覆盖真实算法。63- **切分算法必须抽到 `op_host/<op>_tiling_core.h`**。`<Op>TilingFunc` / `InferShape` 的签名依赖 `gert::TilingContext`,CANN 二进制包未提供 gert faker,无法直接 UT;把与平台无关的算法抽出来后,ophost UT 才能覆盖真实算法。
62- **C++ / CANN 接口**:OpType 用 PascalCase(如 `ComplexMul`),aclnn 为 `aclnn` + OpType(如 `aclnnComplexMul`)。64- **C++ / CANN 接口**:OpType 用 PascalCase(如 `ComplexMul`),aclnn 为 `aclnn` + OpType(如 `aclnnComplexMul`)。
63- **Python 接口**:蛇形命名(如 `complex_mul`),不要用驼峰,也不要加 `npu_` 前缀。65- **Python 接口**:蛇形命名(如 `complex_mul`),不要用驼峰,也不要加 `npu_` 前缀。
64-- 新增算子后,需把它的 `op_host` 目录加进 `tests/ascendc/ophost/CMakeLists.txt` 的 `OPHOST_UT_INCLUDE_DIRS`,并把测试源加进 `ophost_ut` 目标。66+- 新增算子后,在 `tests/ut/op_host/<op>/` 放下 `test_*.cpp`;`tests/ascendc/ophost/CMakeLists.txt` 会自动收集。
67+- 有卡 aclnn:在 `examples/aclnn/<op>/` 放 `test_aclnn_<op>.cpp`(内联小 tensor),在 `tests/st/aclnn<OpType>/` 放 `cases.json` + `gen_data.py` + 两段式 `main.cpp`,并登记 `tests/suites.json`。本仓没有 ATK / `OP_API_UT`,不要空放 `ut/op_api`。
65- pybind 封装:Python API 放在 `ascendc/pybind/<op_family>.py`,C++ 源码放在 `ascendc/pybind/csrc/<op_family>/`,会自动编译成扩展模块 `<op_family>_npu`,不需要改构建脚本。68- pybind 封装:Python API 放在 `ascendc/pybind/<op_family>.py`,C++ 源码放在 `ascendc/pybind/csrc/<op_family>/`,会自动编译成扩展模块 `<op_family>_npu`,不需要改构建脚本。
66 69 
67#### Triton-Ascend 算子(Python whl)70#### Triton-Ascend 算子(Python whl)
@@ -75,7 +78,7 @@ triton-ascend/${op_family}/
75└── README.md # 算子说明78└── README.md # 算子说明
76 79 
77tests/triton/${op_family}/80tests/triton/${op_family}/
78-└── test_triton_${op_family}.py # CPU 回退精度测试(无卡可跑)81+└── test_triton_${op_family}.py # CPU 回退路径精度测试
79```82```
80 83 
81同时需在 `setup.py` 的 `packages` / `package_dir` 中按算子名注册,使其进入 whl84同时需在 `setup.py` 的 `packages` / `package_dir` 中按算子名注册,使其进入 whl
@@ -83,23 +86,21 @@ tests/triton/${op_family}/
83 86 
84约定:87约定:
85 88 
86-- kernel 需提供**非 NPU 环境下的 PyTorch 回退实现**,保证无卡环境可做精度自检。89+- kernel 需提供 **PyTorch 回退实现**,接口语义与 NPU 路径一致。
87- 新增 kernel 分支需能被 `triton-ascend/tools/aot_compile_check.py` 覆盖到。90- 新增 kernel 分支需能被 `triton-ascend/tools/aot_compile_check.py` 覆盖到。
88 91 
89### 4. 精度要求92### 4. 精度要求
90 93 
91-新贡献算子需满足精度标准,具体参见[生态算子开源精度标准](https://gitcode.com/cann/opbase/blob/master/docs/zh/ops_precision_standard/experimental_standard.md)。目录、命名、切分可测与 PR 交付件见[算子开发规格](docs/zh/develop/operator-dev-spec.md)。请在 PR 描述中附上精度测试报告,报告格式参考 [`docs/zh/develop/precision_acceptance_template.md`](docs/zh/develop/precision_acceptance_template.md)。94+新贡献算子需满足精度标准,具体参见[生态算子开源精度标准](https://gitcode.com/cann/opbase/blob/master/docs/zh/ops_precision_standard/experimental_standard.md)。目录、命名、切分可测与 PR 交付件见[算子开发指南](docs/zh/develop/operator_development_guide.md)。请在 PR 描述中附上精度测试报告,报告格式参考 [`docs/zh/develop/precision_acceptance_template.md`](docs/zh/develop/precision_acceptance_template.md)。
92 95 
93-> **算子在 NPU 上的精度、性能、显存必须在真实 910B 上验证,编译通过不等于功能正确。**96+> **算子在 NPU 上的精度、性能、显存必须在真实 910B 上验证通过。**
94 97 
95-### 5. 提交前本地自检(无卡即可执行)98+### 5. 提交前本地自检
96 99 
97```bash100```bash
98-# AscendC:编译打包 OPP .run101+bash build.sh --install
99-bash ascendc/build.sh --pkg102+bash build.sh -u --ut --st --precision --perf
100- 103+bash scripts/check_license_header.sh
101-# 全量无卡单元测试(ophost tiling UT + Triton CPU 回退精度 + Triton AOT 编译)
102-bash build.sh -u
103```104```
104 105 
105### 6. PR 提交106### 6. PR 提交
@@ -108,7 +109,7 @@ bash build.sh -u
108 - C++ 代码符合《[C++ 编程规范](https://gitcode.com/cann/community/blob/master/contributor/coding-standards/C++%20Coding%20standards.md)》,可用仓内 `.clang-format` 格式化。109 - C++ 代码符合《[C++ 编程规范](https://gitcode.com/cann/community/blob/master/contributor/coding-standards/C++%20Coding%20standards.md)》,可用仓内 `.clang-format` 格式化。
109 - 所有新增源文件带 Apache License 2.0 许可证头(参考现有文件)。110 - 所有新增源文件带 Apache License 2.0 许可证头(参考现有文件)。
110 - Markdown 文档语法符合规范。111 - Markdown 文档语法符合规范。
111- - 无卡自检全部通过。112+ - 本地自检全部通过(含 910B 上板)。
112- **PR 提交**:通过 `git` 命令向目标分支提 PR,检查 PR 标题是否清晰、PR 描述是否规范(说明更改内容和原因、关联对应 Issue)、是否已签署 CLA。113- **PR 提交**:通过 `git` 命令向目标分支提 PR,检查 PR 标题是否清晰、PR 描述是否规范(说明更改内容和原因、关联对应 Issue)、是否已签署 CLA。
113 114 
114### 7. CI 门禁115### 7. CI 门禁
@@ -121,13 +122,9 @@ bash build.sh -u
121 122 
122门禁通过后,请在关联的 Issue 中 @ 指派的 Committer。123门禁通过后,请在关联的 Issue 中 @ 指派的 Committer。
123 124 
124-### 8. Committer 检视125+### 8. 代码检视与合入
125 126 
126-Committer 检视后反馈意见,请据此修改,完成后 @ 指派的 Committer。127+合入前需至少两名 Committer 完成代码检视,并分别在 PR 中评论 `/lgtm` 与 `/approve`。检视意见请据此修改,完成后再次 @ 检视人。
127- 
128-### 9. Maintainer 合入
129- 
130-Committer 检视通过后标注 `/lgtm`。Maintainer 将在 1 天内进行最终审核,确认无问题后标注 `/approve` 合入 PR。
131 128 
132---129---
133 130 
@@ -166,7 +163,8 @@ Bug 类 Issue 建议附带:
166- [快速入门](docs/QUICKSTART.md)163- [快速入门](docs/QUICKSTART.md)
167- [目录结构](docs/zh/context/dir_structure.md)164- [目录结构](docs/zh/context/dir_structure.md)
168- [算子列表](docs/zh/op_list.md)165- [算子列表](docs/zh/op_list.md)
169-- [算子开发规格](docs/zh/develop/operator-dev-spec.md)166+- [接口列表](docs/zh/api_list.md)
167+- [算子开发指南](docs/zh/develop/operator_development_guide.md)
170- [安全声明](SECURITY.md)168- [安全声明](SECURITY.md)
171- [许可证](LICENSE)(Apache License 2.0)169- [许可证](LICENSE)(Apache License 2.0)
172- [变更记录](CHANGELOG.md)170- [变更记录](CHANGELOG.md)
MREADME.md+52-28
@@ -2,7 +2,7 @@
2 2 
3## 🚀 概述3## 🚀 概述
4 4 
5-oil-gas-ops-prospect 是面向**油气勘探**(oil & gas exploration)领域的昇腾自定义算子库。仓名即 **oil-gas-ops**(油气算子)+ **prospect**(勘探目标/圈闭),面向地震成像、全波形反演等勘探计算负载,不绑定任何单一上层模型工程。5+oil-gas-ops-prospect 是面向**油气勘探**(oil & gas exploration)领域的昇腾自定义算子库。仓名即 **oil-gas-ops**(油气算子)+ **prospect**(勘探目标),面向地震成像、全波形反演等勘探计算场景。
6 6 
7本仓按**交付物形态**分为两类算子后端;Python 对外接口按算子名暴露,不区分实现后端:7本仓按**交付物形态**分为两类算子后端;Python 对外接口按算子名暴露,不区分实现后端:
8 8 
@@ -15,11 +15,11 @@ oil-gas-ops-prospect 是面向**油气勘探**(oil & gas exploration)领域
15 15 
16当前仓库已验证通过的 CANN 版本如下:16当前仓库已验证通过的 CANN 版本如下:
17 17 
18-| CANN 版本 | 发布时间 | 分支 |18+| CANN 版本 | 发布时间 |
19-|-----------|----------|------|19+|-----------|----------|
20-| [CANN 9.0.0](https://www.hiascend.com/developer/download/community/result?module=cann&cann=9.0.0) | 2026/04/30 | master |20+| [9.0.0 官方版本](https://www.hiascend.com/developer/download/community/result?module=cann&cann=9.0.0) | 2026/04/30 |
21 21 
22-本仓库需要安装 **Toolkit 包**和对应 SoC 的 **Ops 包**,当前面向 **Ascend 910B**(`Ascend910B1`)。完整的安装步骤、SoC 支持矩阵与环境变量配置见[环境部署](docs/zh/context/quick_install.md)。22+本仓库需要安装 **Toolkit 包**、对应 SoC 的 **Ops 包**,以及与 CANN 配套的 **Triton-Ascend**(已验证 `triton-ascend==3.2.1`)。当前面向 **Ascend 910B**(`Ascend910B1`)。完整的安装步骤、SoC 支持矩阵与环境变量配置见[环境部署](docs/zh/context/quick_install.md)。
23 23 
24## ⬇️ 源码下载24## ⬇️ 源码下载
25 25 
@@ -32,12 +32,32 @@ cd oil-gas-ops-prospect
32 32 
33### 编译与安装33### 编译与安装
34 34 
35-在仓库根目录一键完成 OPP 编译 → 安装 → whl 构建 → whl 安装 → 冒烟测试:35+仓库根目录提供三种构建入口。
36+ 
37+**单独编译 OPP `.run` 包**
38+ 
39+```bash
40+bash build.sh --pkg
41+```
42+ 
43+只编译 AscendC OPP `.run`,写入 `output/<SOC>/`,不编 Python whl,也不安装。
44+ 
45+**只编译不安装**
36 46 
37```bash47```bash
38bash build.sh48bash build.sh
39```49```
40 50 
51+同时编译 OPP `.run` 与 Python whl,产物写入 `output/`,不安装到 CANN / Python,也不跑冒烟测试。
52+ 
53+**编译并安装**
54+ 
55+```bash
56+bash build.sh --install
57+```
58+ 
59+编译 OPP 与 whl,安装到 CANN / Python,并跑 whl 冒烟测试。
60+ 
41产物统一落到仓库根目录 `output/`:61产物统一落到仓库根目录 `output/`:
42 62 
43```text63```text
@@ -56,10 +76,11 @@ export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${AS
56 76 
57```bash77```bash
58bash build.sh --help # 查看全部选项78bash build.sh --help # 查看全部选项
79+bash build.sh --pkg -h # Package Options
59bash build.sh -u -h # Test Options80bash build.sh -u -h # Test Options
60-bash build.sh --clean # 先清理再构建81+bash build.sh --clean # 先清理再构建(可与上面组合)
61bash build.sh --ops=complex_mul,complex_mul_grad # 只构建指定算子82bash build.sh --ops=complex_mul,complex_mul_grad # 只构建指定算子
62-bash build.sh --skip-tests # 跳过 whl 冒烟测试83+bash build.sh --install --skip-tests # 安装后跳过 whl 冒烟测试
63bash build.sh --soc=Ascend910B1 -j16 # 指定 SOC 与并行度84bash build.sh --soc=Ascend910B1 -j16 # 指定 SOC 与并行度
64```85```
65 86 
@@ -81,20 +102,19 @@ y = real_irfft(xr, xi, n=256, dim=-1) # 从半谱重建实信号
81 102 
82### 测试103### 测试
83 104 
84-```bash105+分类见 [`tests/README.md`](tests/README.md)。
85-# 全量无卡单元测试:ophost tiling UT + Triton CPU 回退精度 + Triton AOT 编译检查
86-bash build.sh -u
87-# 只跑 ophost / 只编译不执行:
88-# bash build.sh -u --ophost
89-# bash build.sh -u --ophost --noexec
90 106 
91-# 只验证已安装的 whl107+```bash
92-bash tests/run_whl_tests.sh108+bash build.sh -u # 全部套件(需 910B)
109+bash build.sh -u --ut # UT:切分 + 上板小 launch
110+bash build.sh -u --st # ST:多 shape aclnn
111+bash build.sh -u --precision # 精度
112+bash build.sh -u --perf # 性能
113+ 
114+bash tests/run_whl_tests.sh # 只验证已安装的 whl
93```115```
94 116 
95-所有测试用例集中在 [`tests/`](tests/) 下,算子目录只放实现;用例分类与 NPU 依赖见 [`tests/README.md`](tests/README.md)。117+算子在 NPU 上的精度、性能、显存必须在真实 910B 上验证通过。问题定位见[调试与性能分析](docs/zh/debug/op_debug_prof.md)。
96- 
97-> **编译通过不等于功能正确。** 算子在 NPU 上的精度、性能、显存必须在真实 910B 上验证。无卡验证手段与问题定位见[调试与性能分析](docs/zh/debug/op_debug_prof.md)。
98 118 
99## 📖 接口支持清单119## 📖 接口支持清单
100 120 
@@ -122,19 +142,21 @@ oil-gas-ops-prospect/
122│ └── build.sh # AscendC 侧入口(--pkg 打包 / -u --ophost 跑 UT)142│ └── build.sh # AscendC 侧入口(--pkg 打包 / -u --ophost 跑 UT)
123├── triton-ascend/ # Triton-Ascend 算子(real_fft 打包为 oil_gas_ops_prospect.real_fft)143├── triton-ascend/ # Triton-Ascend 算子(real_fft 打包为 oil_gas_ops_prospect.real_fft)
124│ ├── real_fft/ # real_rfft / real_irfft144│ ├── real_fft/ # real_rfft / real_irfft
125-│ └── tools/ # 无卡 AOT 编译检查145+│ └── tools/ # AOT 编译检查
126-├── tests/ # 全部测试用例(算子目录内不放用例)146+├── tests/ # UT / ST / Python 精度性能;operators/ 不再放测试
127-│ ├── ascendc/ophost/ # tiling UT(无卡,googletest)147+│ ├── ut/op_host/ # tiling googletest
148+│ ├── st/aclnn<OpType>/ # 两段式 aclnn 系统测试
149+│ ├── ascendc/ophost/ # 汇总编译 ophost_ut(googletest)
128│ ├── ascendc/complex_mul/# NPU 精度/性能对比150│ ├── ascendc/complex_mul/# NPU 精度/性能对比
129│ ├── ascendc/fused_bias_softmax/151│ ├── ascendc/fused_bias_softmax/
130-│ └── triton/real_fft/ # rFFT / iRFFT 精度(无卡回退 CPU)152+│ └── triton/real_fft/ # rFFT / iRFFT 精度(含 CPU 回退路径)
131-├── examples/ # 可运行示例153+├── examples/ # Python 示例 + examples/aclnn/<op>/ C++ 样例
132├── scripts/ # 辅助脚本154├── scripts/ # 辅助脚本
133├── docs/ # 文档155├── docs/ # 文档
134├── cmake/ # CMake 模块(三方件解析、安全编译选项)156├── cmake/ # CMake 模块(三方件解析、安全编译选项)
135├── .gitcode/workflows/ # GitCode Action 门禁流水线157├── .gitcode/workflows/ # GitCode Action 门禁流水线
136-├── build.sh # 顶层构建入口(-u 跑无卡 UT)158+├── build.sh # 顶层构建入口(默认只编包;--install 安装;-u 跑测试)
137-├── CMakeLists.txt # 顶层 CMake(无卡 C++ 目标:ophost UT)159+├── CMakeLists.txt # 顶层 CMake(ophost tiling UT)
138└── version.cmake # 版本与依赖定义160└── version.cmake # 版本与依赖定义
139```161```
140 162 
@@ -142,12 +164,12 @@ oil-gas-ops-prospect/
142 164 
143## 🛠️ 开发指南165## 🛠️ 开发指南
144 166 
145-新增 AscendC 或 Triton-Ascend 算子的完整流程、目录约定、命名约定与自检方式见[算子开发指南](docs/zh/develop/operator_development_guide.md)与[算子开发规格](docs/zh/develop/operator-dev-spec.md)。167+新增 AscendC 或 Triton-Ascend 算子的完整流程、目录约定、命名约定与自检方式见[算子开发指南](docs/zh/develop/operator_development_guide.md)。
146 168 
147提交前自检:169提交前自检:
148 170 
149```bash171```bash
150-bash build.sh -u # 全量无卡单元测试172+bash build.sh -u # 全部测试套件(需 910B)
151bash scripts/check_license_header.sh # 许可证头检查173bash scripts/check_license_header.sh # 许可证头检查
152```174```
153 175 
@@ -158,6 +180,8 @@ bash scripts/check_license_header.sh # 许可证头检查
158- [目录结构](docs/zh/context/dir_structure.md)180- [目录结构](docs/zh/context/dir_structure.md)
159- [算子列表](docs/zh/op_list.md)181- [算子列表](docs/zh/op_list.md)
160- [接口列表](docs/zh/api_list.md)182- [接口列表](docs/zh/api_list.md)
183+- [编译参数说明](docs/zh/context/build.md)
184+- [算子开发指南](docs/zh/develop/operator_development_guide.md)
161- [贡献指南](CONTRIBUTING.md)185- [贡献指南](CONTRIBUTING.md)
162- [安全声明](SECURITY.md)186- [安全声明](SECURITY.md)
163- [变更记录](CHANGELOG.md)187- [变更记录](CHANGELOG.md)
@@ -37,7 +37,8 @@
37| 三方件声明 | 不涉及 | Third_Party_Open_Source_Software_List.yaml<br>Third_Party_Open_Source_Software_Notice | [https://gitcode.com/cann-src-third-party/googletest](https://gitcode.com/cann-src-third-party/googletest) | googletest 上游获取地址,用于三方件清单声明,不在构建期访问 |37| 三方件声明 | 不涉及 | Third_Party_Open_Source_Software_List.yaml<br>Third_Party_Open_Source_Software_Notice | [https://gitcode.com/cann-src-third-party/googletest](https://gitcode.com/cann-src-third-party/googletest) | googletest 上游获取地址,用于三方件清单声明,不在构建期访问 |
38| 三方件声明 | 不涉及 | Third_Party_Open_Source_Software_List.yaml<br>Third_Party_Open_Source_Software_Notice | [https://github.com/pytorch/pytorch](https://github.com/pytorch/pytorch) | PyTorch 上游获取地址,用于三方件清单声明。PyTorch 由用户自行安装,本仓不下载、不分发 |38| 三方件声明 | 不涉及 | Third_Party_Open_Source_Software_List.yaml<br>Third_Party_Open_Source_Software_Notice | [https://github.com/pytorch/pytorch](https://github.com/pytorch/pytorch) | PyTorch 上游获取地址,用于三方件清单声明。PyTorch 由用户自行安装,本仓不下载、不分发 |
39| 三方件声明 | 不涉及 | Third_Party_Open_Source_Software_List.yaml<br>Third_Party_Open_Source_Software_Notice | [https://gitee.com/ascend/pytorch](https://gitee.com/ascend/pytorch) | torch_npu 上游获取地址,用于三方件清单声明。由用户自行安装,本仓不下载、不分发 |39| 三方件声明 | 不涉及 | Third_Party_Open_Source_Software_List.yaml<br>Third_Party_Open_Source_Software_Notice | [https://gitee.com/ascend/pytorch](https://gitee.com/ascend/pytorch) | torch_npu 上游获取地址,用于三方件清单声明。由用户自行安装,本仓不下载、不分发 |
40-| 三方件声明 | 不涉及 | Third_Party_Open_Source_Software_List.yaml<br>Third_Party_Open_Source_Software_Notice | [https://gitcode.com/cann/triton-ascend](https://gitcode.com/cann/triton-ascend) | Triton-Ascend 上游获取地址,用于三方件清单声明。由用户自行安装,本仓不下载、不分发 |40+| 三方件声明 | 不涉及 | Third_Party_Open_Source_Software_List.yaml<br>Third_Party_Open_Source_Software_Notice | [https://gitcode.com/Ascend/triton-ascend](https://gitcode.com/Ascend/triton-ascend) | Triton-Ascend 上游获取地址,用于三方件清单声明。由用户自行安装,本仓不下载、不分发 |
41+| 依赖 | 不涉及 | docs/zh/context/quick_install.md | [https://mirrors.huaweicloud.com/ascend/repos/pypi](https://mirrors.huaweicloud.com/ascend/repos/pypi) | 用户按环境部署文档 `pip install triton-ascend==3.2.1` 时使用的 extra-index。本仓不下载、不分发 |
41 42 
42---43---
43 44 
@@ -77,7 +77,7 @@ description: >-
77 77 
78## Stage 2:改码78## Stage 2:改码
79 79 
80-一次一个主题。改 tiling 公式必须同步改 `tests/ascendc/ophost/` 手算期望。80+一次一个主题。改 tiling 公式必须同步改对应算子 `tests/ut/op_host/` 手算期望。
81 81 
82改完立刻走 [eval.md](references/eval.md)。精度/运行时失败走 repair-guide。82改完立刻走 [eval.md](references/eval.md)。精度/运行时失败走 repair-guide。
83 83 
@@ -1,6 +1,6 @@
1# 产物1# 产物
2 2 
3-- 算子改动:`ascendc/operators/fused_bias_softmax/`、`ascendc/operators/fused_softmax_grad/`、对应 `tests/ascendc/ophost/`。3+- 算子改动:`ascendc/operators/fused_bias_softmax/`、`ascendc/operators/fused_softmax_grad/`、对应 `tests/ut/op_host/`。
4- 轮次笔记:`extra-info/fused-optimize/`(**不提交**)。4- 轮次笔记:`extra-info/fused-optimize/`(**不提交**)。
5- 每轮至少写:方向、改了什么、ophost / 官方精度、vs Torch 表、keep 或 revert。5- 每轮至少写:方向、改了什么、ophost / 官方精度、vs Torch 表、keep 或 revert。
6- 对外测试报告若更新倍数,只写 Atlas A2 / CANN 9.0.0,不写主机。6- 对外测试报告若更新倍数,只写 Atlas A2 / CANN 9.0.0,不写主机。
@@ -8,7 +8,7 @@
8bash ascendc/build.sh -u --ophost8bash ascendc/build.sh -u --ophost
9```9```
10 10 
11-覆盖 `tests/ascendc/ophost/test_fused_bias_softmax_tiling.cpp` 与 `test_fused_softmax_grad_tiling.cpp`。切分公式一变,先改这些手算期望。11+覆盖 `tests/ut/op_host/fused_bias_softmax/test_fused_bias_softmax_tiling.cpp` 与 `tests/ut/op_host/fused_softmax_grad/test_fused_softmax_grad_tiling.cpp`。切分公式一变,先改这些手算期望。
12 12 
13## NPU 官方精度13## NPU 官方精度
14 14 
@@ -15,7 +15,7 @@ oil-gas-ops-prospect:
15 license: BSD 3-Clause License15 license: BSD 3-Clause License
16 type: run16 type: run
17 triton-ascend:17 triton-ascend:
18- url: https://gitcode.com/cann/triton-ascend18+ url: https://gitcode.com/Ascend/triton-ascend
19 version: matched with CANN19 version: matched with CANN
20 license: MIT License20 license: MIT License
21 type: run21 type: run
@@ -87,7 +87,7 @@ Please refer to https://gitee.com/ascend/pytorch/blob/master/LICENSE for the ful
87 87 
88------------------------------------------------------------------------------------------------88------------------------------------------------------------------------------------------------
89Software: Triton-Ascend89Software: Triton-Ascend
90-Obtained from: https://gitcode.com/cann/triton-ascend90+Obtained from: https://gitcode.com/Ascend/triton-ascend
91Usage: Runtime dependency. Provides the triton.backends.ascend backend used by the Triton91Usage: Runtime dependency. Provides the triton.backends.ascend backend used by the Triton
92 operators under triton-ascend/. Not redistributed with this SOFTWARE.92 operators under triton-ascend/. Not redistributed with this SOFTWARE.
93 93 
@@ -96,7 +96,7 @@ Copyright (c) 2018-2026 Philippe Tillet, OpenAI
96Copyright (c) Huawei Technologies Co., Ltd. 2024-2026. All rights reserved.96Copyright (c) Huawei Technologies Co., Ltd. 2024-2026. All rights reserved.
97 97 
98License: MIT License98License: MIT License
99-Please refer to https://gitcode.com/cann/triton-ascend/blob/master/LICENSE for the full license99+Please refer to https://gitcode.com/Ascend/triton-ascend/blob/master/LICENSE for the full license
100text.100text.
101 101 
102------------------------------------------------------------------------------------------------102------------------------------------------------------------------------------------------------
@@ -44,25 +44,27 @@ output/
44 44 
45## 编译与 UT45## 编译与 UT
46 46 
47-在本目录执行(容器内一般为 `/root/projects/oil-gas-ops-prospect/ascendc`):47+在本目录执行:
48 48 
49```bash49```bash
50# 编译并打包 OPP .run(默认 SOC=Ascend910B1,默认全部算子)50# 编译并打包 OPP .run(默认 SOC=Ascend910B1,默认全部算子)
51# → <repo>/output/Ascend910B1/oil_gas_ops_prospect_opp_<os>_<arch>.run51# → <repo>/output/Ascend910B1/oil_gas_ops_prospect_opp_<os>_<arch>.run
52bash build.sh --pkg52bash build.sh --pkg
53 53 
54-# 运行 ophost 侧单元测试(googletest,无需 NPU 设备)54+# 运行 ophost 侧单元测试(googletest)
55bash build.sh -u --ophost55bash build.sh -u --ophost
56```56```
57 57 
58-测试用例统一收在仓库根目录的 [`tests/`](../test/) 下,本目录只放算子实现。58+C++ 测试放在仓库根 `tests/ut`、`tests/st`;样例在 `examples/aclnn/`。`tests/ascendc/ophost/` 只做汇总编译。`operators/` 只保留 `op_host` / `op_kernel`。
59-ophost UT 工程在 [`tests/ascendc/ophost/`](../tests/ascendc/ophost/),覆盖两部分:
60 59 
61| 测试文件 | 覆盖内容 |60| 测试文件 | 覆盖内容 |
62|----------|----------|61|----------|----------|
63-| `test_complex_mul_tiling.cpp` | 两个算子的 `*_tiling_core.h`:UB 预算约束(前向 13 份 tile / 反向 21 份)、`tileFloats` 的 32B 对齐与上限、小 tensor 收缩、退化输入(负元素数、`aivNum=0`、空张量)、以及 kernel 侧核间切分的覆盖性不变量(每个元素恰好被一个 AIV 核处理一次、区间无空洞无重叠) |62+| `tests/ut/op_host/complex_mul/test_complex_mul_tiling.cpp` | 前向 `*_tiling_core.h`:UB 预算、32B 对齐、小 tensor 收缩、退化输入、核间覆盖性不变量 |
64-| `test_fused_bias_softmax_tiling.cpp` | `FusedBiasSoftmax` 切分:对齐 padding、bias 广播/全尺寸 UB 预算、非法输入拒绝 |63+| `tests/ut/op_host/complex_mul_grad/test_complex_mul_grad_tiling.cpp` | 反向 UB 份数(21)与单拍更短 |
65-| `test_fused_softmax_grad_tiling.cpp` | `FusedSoftmaxGrad` 切分:MIX/VEC 路径分发、V plane 驻留上限、`Q≠S` |64+| `tests/ut/op_host/fused_bias_softmax/test_fused_bias_softmax_tiling.cpp` | 对齐 padding、bias 广播/全尺寸 UB 预算、非法输入拒绝 |
65+| `tests/ut/op_host/fused_softmax_grad/test_fused_softmax_grad_tiling.cpp` | MIX/VEC 路径分发、V plane 驻留上限、`Q≠S` |
66+| `examples/aclnn/<op>/` | 两段式 `aclnn` 小样例(需要 NPU) |
67+| `tests/st/aclnn<OpType>/` | 多 case 两段式 `aclnn` + 真值对比(需要 NPU) |
66 68 
67> `<Op>TilingFunc` / `InferShape` 的入口签名依赖 `gert::TilingContext`,`Compute<Op>Tiling`69> `<Op>TilingFunc` / `InferShape` 的入口签名依赖 `gert::TilingContext`,`Compute<Op>Tiling`
68> 依赖 `platform_ascendc::PlatformAscendC`,CANN 二进制包未提供 gert faker(该基建在70> 依赖 `platform_ascendc::PlatformAscendC`,CANN 二进制包未提供 gert faker(该基建在
@@ -70,8 +72,7 @@ ophost UT 工程在 [`tests/ascendc/ophost/`](../tests/ascendc/ophost/),覆盖
70> `op_host/<op>_tiling_core.h`,UT 直接驱动核心函数,覆盖的是真实算法。产物在72> `op_host/<op>_tiling_core.h`,UT 直接驱动核心函数,覆盖的是真实算法。产物在
71> `_work/ophost_ut/`。73> `_work/ophost_ut/`。
72>74>
73-> 新增算子时:把它的 `op_host` 目录加进 `tests/ascendc/ophost/CMakeLists.txt` 的75+> 新增算子时:在 `tests/ut/op_host/<op>/` 放下 `test_*.cpp`,CMake 会自动收集。
74-> `OPHOST_UT_INCLUDE_DIRS`,并把测试源加进 `ophost_ut` 目标。
75 76 
76常用打包参数:77常用打包参数:
77 78 
@@ -90,13 +91,12 @@ bash build.sh list-ops
90python3 build.py -v Ascend910B1 --ops complex_mul complex_mul_grad91python3 build.py -v Ascend910B1 --ops complex_mul complex_mul_grad
91```92```
92 93 
93-## 一键构建并安装 OPP + whl94+## 编译与安装 OPP + whl
94 95 
95如果目标是直接使用自定义算子(含 whl),优先在仓库根目录执行:96如果目标是直接使用自定义算子(含 whl),优先在仓库根目录执行:
96 97 
97```bash98```bash
98-cd /root/projects/oil-gas-ops-prospect99+bash build.sh --install
99-bash build.sh
100```100```
101 101 
102该脚本会自动完成:102该脚本会自动完成:
@@ -110,18 +110,21 @@ bash build.sh
110常用参数:110常用参数:
111 111 
112```bash112```bash
113+bash build.sh --pkg # 只编译 OPP .run
114+bash build.sh # 编译 .run 与 whl,不安装
115+bash build.sh --install # 编译并安装
113bash build.sh --clean # 先清理 ascendc/_work 与 output/116bash build.sh --clean # 先清理 ascendc/_work 与 output/
114bash build.sh --ops=complex_mul # 只构建指定 AscendC 算子117bash build.sh --ops=complex_mul # 只构建指定 AscendC 算子
115-bash build.sh --skip-tests # 跳过 whl 冒烟测试118+bash build.sh --install --skip-tests # 安装后跳过 whl 冒烟测试
116bash build.sh --soc=Ascend910B1 -j16119bash build.sh --soc=Ascend910B1 -j16
117-bash build.sh -u # 无卡 UT120+bash build.sh -u # 全部测试套件(需 910B)
118bash build.sh -u --ophost --noexec # 只编译 ophost UT121bash build.sh -u --ophost --noexec # 只编译 ophost UT
119```122```
120 123 
121如 CANN 安装路径不是默认的 `/usr/local/Ascend/cann-9.0.0`,可通过环境变量指定:124如 CANN 安装路径不是默认的 `/usr/local/Ascend/cann-9.0.0`,可通过环境变量指定:
122 125 
123```bash126```bash
124-ASCEND_HOME_PATH=/path/to/cann bash build.sh127+ASCEND_HOME_PATH=/path/to/cann bash build.sh --install
125```128```
126 129 
127### 包名与 vendor130### 包名与 vendor
@@ -184,7 +187,7 @@ CANN 内部框架头(`op_host/tiling_base.h`、`op_host/tiling_templates_regis
184 187 
185| 目录名 | OpType | aclnn 接口 | Python | 说明 |188| 目录名 | OpType | aclnn 接口 | Python | 说明 |
186|--------|--------|-----------|--------|------|189|--------|--------|-----------|--------|------|
187-| `complex_mul` | `ComplexMul` | `aclnnComplexMul` | `oil_gas_ops_prospect.complex_mul.complex_mul` | 复数逐点乘(AIV-only) |190+| `complex_mul` | `ComplexMul` | `aclnnComplexMul` | `oil_gas_ops_prospect.complex_mul.complex_mul` | 四路 float 复数逐点乘(AIV-only) |
188| `complex_mul_grad` | `ComplexMulGrad` | `aclnnComplexMulGrad` | 由 `complex_mul` 的 autograd 调用 | 上者的反向(AIV-only) |191| `complex_mul_grad` | `ComplexMulGrad` | `aclnnComplexMulGrad` | 由 `complex_mul` 的 autograd 调用 | 上者的反向(AIV-only) |
189| `fused_bias_softmax` | `FusedBiasSoftmax` | `aclnnFusedBiasSoftmax` | `oil_gas_ops_prospect.fused_bias_softmax.fused_bias_softmax` | `softmax(x+bias1+bias2,-1)`,原地写回 |192| `fused_bias_softmax` | `FusedBiasSoftmax` | `aclnnFusedBiasSoftmax` | `oil_gas_ops_prospect.fused_bias_softmax.fused_bias_softmax` | `softmax(x+bias1+bias2,-1)`,原地写回 |
190| `fused_softmax_grad` | `FusedSoftmaxGrad` | `aclnnFusedSoftmaxGrad` | `oil_gas_ops_prospect.fused_bias_softmax.fused_softmax_grad` | 注意力 softmax JVP,原地覆写 probs |193| `fused_softmax_grad` | `FusedSoftmaxGrad` | `aclnnFusedSoftmaxGrad` | `oil_gas_ops_prospect.fused_bias_softmax.fused_softmax_grad` | 注意力 softmax JVP,原地覆写 probs |
@@ -247,7 +250,8 @@ Python 还提供 `attention_core(q, k, v, bias1, bias2)`,把上述两算子按
247> C++ 侧 OpType 为 PascalCase(`ComplexMul`),aclnn 为 `aclnnComplexMul`。250> C++ 侧 OpType 为 PascalCase(`ComplexMul`),aclnn 为 `aclnnComplexMul`。
248> Python 侧公开函数为蛇形命名(`complex_mul`),不加 `npu_` 前缀。251> Python 侧公开函数为蛇形命名(`complex_mul`),不加 `npu_` 前缀。
249> kernel 入口通过 `opFile.value`/`opInterface.value` 固定为蛇形文件名252> kernel 入口通过 `opFile.value`/`opInterface.value` 固定为蛇形文件名
250-> (`complex_mul` / `complex_mul_grad` / `fused_bias_softmax` / `fused_softmax_grad`),253+> (`complex_mul` / `complex_mul_grad` /
254+> `fused_bias_softmax` / `fused_softmax_grad`),
251> OpType 改名不影响 kernel 源码;aclnn 接口随 OpType 变化。255> OpType 改名不影响 kernel 源码;aclnn 接口随 OpType 变化。
252 256 
253## 测试用例257## 测试用例
@@ -270,13 +274,12 @@ tests/ascendc/<op>/
270测试用到的 pybind 扩展是算子交付件而非测试资产,源码在274测试用到的 pybind 扩展是算子交付件而非测试资产,源码在
271[`pybind/`](pybind/),详见 [`pybind/README.md`](pybind/README.md)。275[`pybind/`](pybind/),详见 [`pybind/README.md`](pybind/README.md)。
272 276 
273-另有不依赖 NPU 的 ophost UT,见上文「编译与 UT」。277+另有 ophost tiling UT,见上文「编译与 UT」。
274 278 
275### 前置条件279### 前置条件
276 280 
277-1. 已在 `ascendc-dev` 容器中构建出 OPP 包(见上文「编译与 UT」),产物在 `<repo>/output/<SOC>/oil_gas_ops_prospect_opp_*.run`。281+1. 已构建出 OPP 包(见上文「编译与 UT」),产物在 `<repo>/output/<SOC>/oil_gas_ops_prospect_opp_*.run`。
2782. **测试需在带 NPU 设备的机器上运行**,且环境内有 `torch` 与 `torch_npu`。2822. **测试需在带 NPU 设备的机器上运行**,且环境内有 `torch` 与 `torch_npu`。
279- 编译可在容器内完成,跑测试需真实 NPU。
280 283 
281### 运行284### 运行
282 285 
@@ -0,0 +1,76 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * Licensed under the Apache License, Version 2.0 (the "License");
4+ * you may not use this file except in compliance with the License.
5+ * You may obtain a copy of the License at
6+ *
7+ * http://www.apache.org/licenses/LICENSE-2.0
8+ *
9+ * Unless required by applicable law or agreed to in writing, software
10+ * distributed under the License is distributed on an "AS IS" BASIS,
11+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+ * See the License for the specific language governing permissions and
13+ * limitations under the License.
14+ */
15+ 
16+/**
17+ * 本仓自定义算子的 aclnn 两段式接口声明,供 examples/ 与 tests/st/ 共用。
18+ *
19+ * 装了算子包时优先用 op_api/include 下生成的 aclnn_*.h;没装时回落到手写声明,
20+ * 这样源码树里的 examples / ST 工程仍可独立编译。声明按算子逐个用 __has_include
21+ * 判定,只装了部分算子也能正常编译。
22+ */
23+ 
24+#ifndef OG_ACLNN_DECLS_H
25+#define OG_ACLNN_DECLS_H
26+ 
27+#include "acl/acl.h"
28+#include "aclnn/acl_meta.h"
29+ 
30+#if __has_include("aclnn_complex_mul.h")
31+#include "aclnn_complex_mul.h"
32+#else
33+extern "C" {
34+aclnnStatus aclnnComplexMulGetWorkspaceSize(const aclTensor* ur, const aclTensor* ui, const aclTensor* vr,
35+ const aclTensor* vi, const aclTensor* wrOut, const aclTensor* wiOut,
36+ uint64_t* workspaceSize, aclOpExecutor** executor);
37+aclnnStatus aclnnComplexMul(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
38+}
39+#endif
40+ 
41+#if __has_include("aclnn_complex_mul_grad.h")
42+#include "aclnn_complex_mul_grad.h"
43+#else
44+extern "C" {
45+aclnnStatus aclnnComplexMulGradGetWorkspaceSize(const aclTensor* gwr, const aclTensor* gwi, const aclTensor* ur,
46+ const aclTensor* ui, const aclTensor* vr, const aclTensor* vi,
47+ const aclTensor* gurOut, const aclTensor* guiOut,
48+ const aclTensor* gvrOut, const aclTensor* gviOut,
49+ uint64_t* workspaceSize, aclOpExecutor** executor);
50+aclnnStatus aclnnComplexMulGrad(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
51+}
52+#endif
53+ 
54+#if __has_include("aclnn_fused_bias_softmax.h")
55+#include "aclnn_fused_bias_softmax.h"
56+#else
57+extern "C" {
58+aclnnStatus aclnnFusedBiasSoftmaxGetWorkspaceSize(const aclTensor* x, const aclTensor* bias1, const aclTensor* bias2,
59+ const aclTensor* y, uint64_t* workspaceSize,
60+ aclOpExecutor** executor);
61+aclnnStatus aclnnFusedBiasSoftmax(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
62+}
63+#endif
64+ 
65+#if __has_include("aclnn_fused_softmax_grad.h")
66+#include "aclnn_fused_softmax_grad.h"
67+#else
68+extern "C" {
69+aclnnStatus aclnnFusedSoftmaxGradGetWorkspaceSize(const aclTensor* probs, const aclTensor* gradOutput,
70+ const aclTensor* values, const aclTensor* out,
71+ uint64_t* workspaceSize, aclOpExecutor** executor);
72+aclnnStatus aclnnFusedSoftmaxGrad(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
73+}
74+#endif
75+ 
76+#endif // OG_ACLNN_DECLS_H
@@ -45,21 +45,17 @@ bool CheckBiasValid(gert::TilingContext* context, uint32_t idx, const gert::Shap
45 if (shapeVal.GetDimNum() != OgOpsCommon::kOgDimNum || shapeVal.GetShapeSize() == 0) {45 if (shapeVal.GetDimNum() != OgOpsCommon::kOgDimNum || shapeVal.GetShapeSize() == 0) {
46 return false;46 return false;
47 }47 }
48- if (shapeVal.GetDim(OgOpsCommon::kOgDim0) != xShape.GetDim(OgOpsCommon::kOgDim0) ||48+ if (!FusedBiasSoftmaxHost::BiasShapeValid(static_cast<int64_t>(xShape.GetDim(OgOpsCommon::kOgDim0)),
49- shapeVal.GetDim(OgOpsCommon::kOgDim1) != xShape.GetDim(OgOpsCommon::kOgDim1) ||49+ static_cast<int64_t>(xShape.GetDim(OgOpsCommon::kOgDim1)),
50- shapeVal.GetDim(OgOpsCommon::kOgDim3) != xShape.GetDim(OgOpsCommon::kOgDim3)) {50+ static_cast<int64_t>(xShape.GetDim(OgOpsCommon::kOgDim2)),
51+ static_cast<int64_t>(xShape.GetDim(OgOpsCommon::kOgDim3)),
52+ static_cast<int64_t>(shapeVal.GetDim(OgOpsCommon::kOgDim0)),
53+ static_cast<int64_t>(shapeVal.GetDim(OgOpsCommon::kOgDim1)),
54+ static_cast<int64_t>(shapeVal.GetDim(OgOpsCommon::kOgDim2)),
55+ static_cast<int64_t>(shapeVal.GetDim(OgOpsCommon::kOgDim3)),
56+ b1FullRequired == 1)) {
51 return false;57 return false;
52 }58 }
53- if (b1FullRequired == 1) {
54- if (shapeVal.GetDim(OgOpsCommon::kOgDim2) != xShape.GetDim(OgOpsCommon::kOgDim2)) {
55- return false;
56- }
57- } else {
58- if (shapeVal.GetDim(OgOpsCommon::kOgDim2) != 1 &&
59- shapeVal.GetDim(OgOpsCommon::kOgDim2) != xShape.GetDim(OgOpsCommon::kOgDim2)) {
60- return false;
61- }
62- }
63 return desc->GetDataType() == context->GetInputDesc(kShapeInput)->GetDataType();59 return desc->GetDataType() == context->GetInputDesc(kShapeInput)->GetDataType();
64}60}
65 61 
@@ -67,6 +67,25 @@ struct FusedBiasSoftmaxPlan {
67 67 
68inline uint64_t FbsCeilDiv(uint64_t a, uint64_t b) { return (b == 0) ? 0 : (a + b - 1) / b; }68inline uint64_t FbsCeilDiv(uint64_t a, uint64_t b) { return (b == 0) ? 0 : (a + b - 1) / b; }
69 69 
70+/**
71+ * bias 4D 与 x 的 B/H/S 对齐;fullRequired 时 Q 必须等于 xQ,否则 Q 为 1(广播)或 xQ。
72+ * 不依赖 gert,供 tiling 入口与 ophost UT 共用。
73+ */
74+inline bool BiasShapeValid(int64_t xB, int64_t xH, int64_t xQ, int64_t xS, int64_t bB, int64_t bH, int64_t bQ,
75+ int64_t bS, bool fullRequired)
76+{
77+ if (xB <= 0 || xH <= 0 || xQ <= 0 || xS <= 0 || bB <= 0 || bH <= 0 || bQ <= 0 || bS <= 0) {
78+ return false;
79+ }
80+ if (bB != xB || bH != xH || bS != xS) {
81+ return false;
82+ }
83+ if (fullRequired) {
84+ return bQ == xQ;
85+ }
86+ return bQ == 1 || bQ == xQ;
87+}
88+ 
70/** width 对齐到 32B 的元素数与 padding 个数。 */89/** width 对齐到 32B 的元素数与 padding 个数。 */
71inline void ComputePadding(uint64_t width, uint64_t dtypeSize, uint64_t& paddingNum, uint64_t& padLineNum)90inline void ComputePadding(uint64_t width, uint64_t dtypeSize, uint64_t& paddingNum, uint64_t& padLineNum)
72{91{
@@ -3,7 +3,7 @@
3各 AscendC 算子的 **C++ 绑定** 与 **Python 对外接口**(autograd 接线、torch 回退、实现选择)。3各 AscendC 算子的 **C++ 绑定** 与 **Python 对外接口**(autograd 接线、torch 回退、实现选择)。
4绑定层直接调用 msopgen 生成的 `aclnn` 接口(如 `aclnnComplexMul`)。4绑定层直接调用 msopgen 生成的 `aclnn` 接口(如 `aclnnComplexMul`)。
5 5 
6-测试用例统一在仓库根目录的 [`tests/`](../../test/) 下。6+测试用例统一在仓库根目录的 [`tests/`](../../tests/) 下。
7 7 
8## 目录结构8## 目录结构
9 9 
@@ -36,7 +36,7 @@ Python API 放在 `ascendc/pybind/<op_family>.py`(与 `csrc/<op_family>/` 并
36正式形态是随 whl 一起构建(仓库根目录执行):36正式形态是随 whl 一起构建(仓库根目录执行):
37 37 
38```bash38```bash
39-bash build.sh # 构建 OPP 包 → 安装 → 构建并安装 whl(含本扩展)39+bash build.sh --install # 构建 OPP 包 → 安装 → 构建并安装 whl(含本扩展)
40```40```
41 41 
42免装 whl 的就地编译,产物 `<op_family>_npu*.so` 落在本目录:42免装 whl 的就地编译,产物 `<op_family>_npu*.so` 落在本目录:
@@ -51,6 +51,6 @@ python3 ascendc/pybind_setup.py build_ext --inplace
51 51 
52- **OPP 包已安装到 CANN**,即 `opp/vendors/customize/op_api/lib/libcust_opapi.so` 存在。52- **OPP 包已安装到 CANN**,即 `opp/vendors/customize/op_api/lib/libcust_opapi.so` 存在。
53- 环境内有 `torch` 与 `torch_npu`。53- 环境内有 `torch` 与 `torch_npu`。
54-- 编译可在无卡容器内完成,**运行需要真实 NPU**。54+- 运行需要真实 NPU。
55 55 
56置 `OIL_GAS_OPS_SKIP_EXT=1` 可跳过扩展编译,只构建纯 Python 部分。56置 `OIL_GAS_OPS_SKIP_EXT=1` 可跳过扩展编译,只构建纯 Python 部分。
@@ -74,23 +74,34 @@ def torch_complex_mul_grad(
74 return gur, gui, gvr, gvi74 return gur, gui, gvr, gvi
75 75 
76 76 
77+def _kernel_in(t: torch.Tensor, *, clone: bool = False) -> torch.Tensor:
78+ t = t.detach().contiguous()
79+ return t.clone() if clone else t
80+ 
81+ 
77def _ascendc_forward(ur, ui, vr, vi):82def _ascendc_forward(ur, ui, vr, vi):
78 ext = _load_extension()83 ext = _load_extension()
79- return ext.complex_mul_forward(84+ wr, wi = ext.complex_mul_forward(
80- ur.contiguous(), ui.contiguous(), vr.contiguous(), vi.contiguous()85+ _kernel_in(ur),
86+ _kernel_in(ui),
87+ _kernel_in(vr),
88+ _kernel_in(vi),
81 )89 )
90+ # pybind 把 std::vector 变成 list;autograd.Function 多输出必须是独立 tensor 的 tuple。
91+ return wr.clone(), wi.clone()
82 92 
83 93 
84def _ascendc_backward(gwr, gwi, ur, ui, vr, vi):94def _ascendc_backward(gwr, gwi, ur, ui, vr, vi):
85 ext = _load_extension()95 ext = _load_extension()
86- return ext.complex_mul_backward(96+ gur, gui, gvr, gvi = ext.complex_mul_backward(
87- gwr.contiguous(),97+ _kernel_in(gwr, clone=True),
88- gwi.contiguous(),98+ _kernel_in(gwi, clone=True),
89- ur.contiguous(),99+ _kernel_in(ur),
90- ui.contiguous(),100+ _kernel_in(ui),
91- vr.contiguous(),101+ _kernel_in(vr),
92- vi.contiguous(),102+ _kernel_in(vi),
93 )103 )
104+ return gur, gui, gvr, gvi
94 105 
95 106 
96class _ComplexMulFn(torch.autograd.Function):107class _ComplexMulFn(torch.autograd.Function):
@@ -40,7 +40,8 @@ aclnnStatus aclnnComplexMulGrad(void* workspace, uint64_t workspaceSize, aclOpEx
40 40 
41namespace {41namespace {
42 42 
43-using OgPybind::BindTensorAddr;43+using OgPybind::BindInputAddr;
44+using OgPybind::BindOutputAddr;
44using OgPybind::Destroy;45using OgPybind::Destroy;
45using OgPybind::MakeExecutorRepeatable;46using OgPybind::MakeExecutorRepeatable;
46using OgPybind::ShapeOf;47using OgPybind::ShapeOf;
@@ -86,8 +87,9 @@ std::vector<at::Tensor> complex_mul_forward(const at::Tensor& ur, const at::Tens
86 87 
87 const auto shape = ShapeOf(ur);88 const auto shape = ShapeOf(ur);
88 const int64_t numel = ur.numel();89 const int64_t numel = ur.numel();
89- // 每次调用独立输出。TBE 只认第一路基址,wr/wi 必须落在同一块连续缓冲。90+ // TBE 只认第一路基址,wr/wi 必须落在同一块连续缓冲。
90- at::Tensor packed = at::empty({2, numel}, ur.options());91+ // 不要继承输入的 requires_grad,否则 autograd.Function 无法给输出挂自己的 grad_fn。
92+ at::Tensor packed = at::empty({2, numel}, ur.options().requires_grad(false));
91 at::Tensor wr = packed.select(0, 0).view(ur.sizes());93 at::Tensor wr = packed.select(0, 0).view(ur.sizes());
92 at::Tensor wi = packed.select(0, 1).view(ur.sizes());94 at::Tensor wi = packed.select(0, 1).view(ur.sizes());
93 95 
@@ -104,12 +106,12 @@ std::vector<at::Tensor> complex_mul_forward(const at::Tensor& ur, const at::Tens
104 auto ret = aclnnComplexMulGetWorkspaceSize(urT, uiT, vrT, viT, wrT, wiT, &workspaceSize, &executor);106 auto ret = aclnnComplexMulGetWorkspaceSize(urT, uiT, vrT, viT, wrT, wiT, &workspaceSize, &executor);
105 TORCH_CHECK(ret == ACL_SUCCESS, "aclnnComplexMulGetWorkspaceSize failed, ret=", static_cast<int>(ret));107 TORCH_CHECK(ret == ACL_SUCCESS, "aclnnComplexMulGetWorkspaceSize failed, ret=", static_cast<int>(ret));
106 MakeExecutorRepeatable(executor);108 MakeExecutorRepeatable(executor);
107- BindTensorAddr(executor, 0, urT, ur.data_ptr());109+ BindInputAddr(executor, 0, urT, ur.data_ptr());
108- BindTensorAddr(executor, 1, uiT, ui.data_ptr());110+ BindInputAddr(executor, 1, uiT, ui.data_ptr());
109- BindTensorAddr(executor, 2, vrT, vr.data_ptr());111+ BindInputAddr(executor, 2, vrT, vr.data_ptr());
110- BindTensorAddr(executor, 3, viT, vi.data_ptr());112+ BindInputAddr(executor, 3, viT, vi.data_ptr());
111- BindTensorAddr(executor, 4, wrT, packed.data_ptr());113+ BindOutputAddr(executor, 0, wrT, packed.data_ptr());
112- BindTensorAddr(executor, 5, wiT, packed.data_ptr());114+ BindOutputAddr(executor, 1, wiT, packed.data_ptr());
113 115 
114 void* workspace = nullptr;116 void* workspace = nullptr;
115 at::Tensor& workspaceTensor = StableWorkspace(false, workspaceSize, ur);117 at::Tensor& workspaceTensor = StableWorkspace(false, workspaceSize, ur);
@@ -153,16 +155,16 @@ void LaunchComplexMulGrad(const at::Tensor& gwr, const at::Tensor& gwi, const at
153 &workspaceSize, &executor);155 &workspaceSize, &executor);
154 TORCH_CHECK(ret == ACL_SUCCESS, "aclnnComplexMulGradGetWorkspaceSize failed, ret=", static_cast<int>(ret));156 TORCH_CHECK(ret == ACL_SUCCESS, "aclnnComplexMulGradGetWorkspaceSize failed, ret=", static_cast<int>(ret));
155 MakeExecutorRepeatable(executor);157 MakeExecutorRepeatable(executor);
156- BindTensorAddr(executor, 0, gwrT, gwr.data_ptr());158+ BindInputAddr(executor, 0, gwrT, gwr.data_ptr());
157- BindTensorAddr(executor, 1, gwiT, gwi.data_ptr());159+ BindInputAddr(executor, 1, gwiT, gwi.data_ptr());
158- BindTensorAddr(executor, 2, urT, ur.data_ptr());160+ BindInputAddr(executor, 2, urT, ur.data_ptr());
159- BindTensorAddr(executor, 3, uiT, ui.data_ptr());161+ BindInputAddr(executor, 3, uiT, ui.data_ptr());
160- BindTensorAddr(executor, 4, vrT, vr.data_ptr());162+ BindInputAddr(executor, 4, vrT, vr.data_ptr());
161- BindTensorAddr(executor, 5, viT, vi.data_ptr());163+ BindInputAddr(executor, 5, viT, vi.data_ptr());
162- BindTensorAddr(executor, 6, gurT, packed.data_ptr());164+ BindOutputAddr(executor, 0, gurT, packed.data_ptr());
163- BindTensorAddr(executor, 7, guiT, packed.data_ptr());165+ BindOutputAddr(executor, 1, guiT, packed.data_ptr());
164- BindTensorAddr(executor, 8, gvrT, packed.data_ptr());166+ BindOutputAddr(executor, 2, gvrT, packed.data_ptr());
165- BindTensorAddr(executor, 9, gviT, packed.data_ptr());167+ BindOutputAddr(executor, 3, gviT, packed.data_ptr());
166 168 
167 void* workspace = nullptr;169 void* workspace = nullptr;
168 at::Tensor& workspaceTensor = StableWorkspace(true, workspaceSize, ur);170 at::Tensor& workspaceTensor = StableWorkspace(true, workspaceSize, ur);
@@ -189,7 +191,7 @@ std::vector<at::Tensor> complex_mul_backward(const at::Tensor& gwr, const at::Te
189 TORCH_CHECK(gwr.scalar_type() == at::kFloat, "complex_mul_grad requires float32");191 TORCH_CHECK(gwr.scalar_type() == at::kFloat, "complex_mul_grad requires float32");
190 192 
191 const int64_t numel = ur.numel();193 const int64_t numel = ur.numel();
192- at::Tensor packed = at::empty({4, numel}, ur.options());194+ at::Tensor packed = at::empty({4, numel}, ur.options().requires_grad(false));
193 at::Tensor gur = packed.select(0, 0).view(ur.sizes());195 at::Tensor gur = packed.select(0, 0).view(ur.sizes());
194 at::Tensor gui = packed.select(0, 1).view(ur.sizes());196 at::Tensor gui = packed.select(0, 1).view(ur.sizes());
195 at::Tensor gvr = packed.select(0, 2).view(ur.sizes());197 at::Tensor gvr = packed.select(0, 2).view(ur.sizes());
Mbuild.sh+155-122
@@ -25,87 +25,59 @@ WHEEL_DIR="${OUTPUT_DIR}/wheels"
25# shellcheck source=scripts/build_options.sh25# shellcheck source=scripts/build_options.sh
26source "${SCRIPT_DIR}/scripts/build_options.sh"26source "${SCRIPT_DIR}/scripts/build_options.sh"
27 27 
28-run_ophost_ut() {28+run_unified_tests() {
29- local build_dir="${ASCENDC_ROOT}/_work/ophost_ut"29+ local suites=()
30- local ut_src="${SCRIPT_DIR}/tests/ascendc/ophost"30+ [[ "${OP_HOST_UT}" == "TRUE" ]] && suites+=("ophost")
31- local cmake_flags=(-DCMAKE_BUILD_TYPE=Release)31+ [[ "${TRITON_UT}" == "TRUE" ]] && suites+=("triton")
32+ [[ "${AOT_UT}" == "TRUE" ]] && suites+=("aot")
33+ [[ "${ST_UT}" == "TRUE" ]] && suites+=("st")
34+ [[ "${PRECISION_UT}" == "TRUE" ]] && suites+=("precision")
35+ [[ "${PERF_UT}" == "TRUE" ]] && suites+=("perf")
36+ [[ "${EXAMPLES_UT}" == "TRUE" ]] && suites+=("examples")
37+ [[ "${FFT_UT}" == "TRUE" ]] && suites+=("fft")
38+ [[ "${GATE_UT}" == "TRUE" ]] && suites+=("gate")
32 39 
33- if [[ "${ENABLE_COVERAGE}" == "TRUE" ]]; then40+ if [[ ${#suites[@]} -eq 0 ]]; then
34- cmake_flags+=(-DCMAKE_CXX_FLAGS=--coverage -DCMAKE_EXE_LINKER_FLAGS=--coverage)41+ echo "[ERROR] no test suites selected" >&2
42+ exit 1
35 fi43 fi
36 44 
37- echo "[ut] ophost tiling UT (googletest)..."45+ local suite_csv
38- cmake -S "${ut_src}" -B "${build_dir}" "${cmake_flags[@]}"46+ suite_csv="$(IFS=,; echo "${suites[*]}")"
39- cmake --build "${build_dir}" -j "${THREAD_NUM}"
40- 
41- if [[ "${ENABLE_UT_EXEC}" != "TRUE" ]]; then
42- echo "[ut] --noexec: skip running ophost_ut"
43- return 0
44- fi
45 47 
48+ # 取值一律用 --flag=value。裸 --orig-argv -u 会被 argparse 当成新选项,
49+ # 导致 `bash build.sh -u` / `--ut` / `--st` 等单 flag 入口立刻失败。
50+ local args=(
51+ python3 -u "${SCRIPT_DIR}/scripts/run_tests.py"
52+ --suites="${suite_csv}"
53+ --jobs="${THREAD_NUM}"
54+ --orig-argv="${USE_CMD}"
55+ --results-json="${SCRIPT_DIR}/test_result/test_results.json"
56+ )
46 if [[ -n "${GTEST_FILTER}" ]]; then57 if [[ -n "${GTEST_FILTER}" ]]; then
47- echo "[ut] --gtest_filter=${GTEST_FILTER}"58+ args+=(--gtest-filter="${GTEST_FILTER}")
48- "${build_dir}/ophost_ut" --gtest_filter="${GTEST_FILTER}"
49- else
50- ctest --test-dir "${build_dir}" --output-on-failure
51 fi59 fi
52-}
53- 
54-run_triton_ut() {
55- local pytest_args=("${SCRIPT_DIR}/tests/triton/real_fft" -q)
56- 
57 if [[ "${ENABLE_UT_EXEC}" != "TRUE" ]]; then60 if [[ "${ENABLE_UT_EXEC}" != "TRUE" ]]; then
58- echo "[ut] --noexec: skip Triton CPU-fallback pytest"61+ args+=(--noexec)
59- return 0
60 fi62 fi
61- 
62 if [[ "${ENABLE_COVERAGE}" == "TRUE" ]]; then63 if [[ "${ENABLE_COVERAGE}" == "TRUE" ]]; then
63- if python3 -c "import pytest_cov" >/dev/null 2>&1; then64+ args+=(--cov)
64- pytest_args+=(--cov --cov-report=term-missing)65+ fi
65- else66+ if [[ "${REQUIRE_NPU}" == "TRUE" ]]; then
66- echo "[WARN] --cov requested but pytest-cov is not installed; skip Python coverage"67+ args+=(--require-npu)
67- fi68+ fi
69+ if [[ -n "${REPORT_PATH}" ]]; then
70+ args+=(--report="${REPORT_PATH}")
71+ fi
72+ if [[ -n "${OPENSOURCE_REPORT}" ]]; then
73+ args+=(--opensource-report="${OPENSOURCE_REPORT}")
74+ fi
75+ if [[ "${OPS_LIST}" != "all" && -n "${OPS_LIST}" ]]; then
76+ args+=(--ops="${OPS_LIST}")
68 fi77 fi
69 78 
70- echo "[ut] Triton CPU-fallback accuracy tests..."79+ echo "[oil_gas_ops_prospect] test suites: ${suite_csv}"
71- python3 -m pytest "${pytest_args[@]}"80+ "${args[@]}"
72-}
73- 
74-run_aot_check() {
75- if [[ "${ENABLE_UT_EXEC}" != "TRUE" ]]; then
76- echo "[ut] --noexec: skip Triton AOT compile check"
77- return 0
78- fi
79- echo "[ut] Triton AOT compile check..."
80- python3 "${SCRIPT_DIR}/triton-ascend/tools/aot_compile_check.py"
81-}
82- 
83-run_device_free_ut() {
84- local run_ophost=FALSE
85- local run_triton=FALSE
86- local run_aot=FALSE
87- 
88- if [[ "${UT_TEST_ALL}" == "TRUE" || "${OP_HOST_UT}" == "TRUE" ]]; then
89- run_ophost=TRUE
90- fi
91- if [[ "${UT_TEST_ALL}" == "TRUE" || "${TRITON_UT}" == "TRUE" ]]; then
92- run_triton=TRUE
93- fi
94- if [[ "${UT_TEST_ALL}" == "TRUE" || "${AOT_UT}" == "TRUE" ]]; then
95- run_aot=TRUE
96- fi
97- 
98- if [[ "${run_ophost}" == "TRUE" ]]; then
99- run_ophost_ut
100- fi
101- if [[ "${run_triton}" == "TRUE" ]]; then
102- run_triton_ut
103- fi
104- if [[ "${run_aot}" == "TRUE" ]]; then
105- run_aot_check
106- fi
107- 
108- echo "[oil_gas_ops_prospect] Device-free unit tests passed."
109}81}
110 82 
111split_ops_list() {83split_ops_list() {
@@ -117,18 +89,35 @@ for item in raw.split():
117PY89PY
118}90}
119 91 
120-run_full_pipeline() {92+latest_artifact() {
121- echo "[oil_gas_ops_prospect] SOC_VERSION=${SOC_VERSION}"93+ python3 - <<'PY' "$1" "$2"
122- echo "[oil_gas_ops_prospect] OPS_LIST=${OPS_LIST}"94+from pathlib import Path
95+import sys
123 96 
97+out_dir = Path(sys.argv[1])
98+pattern = sys.argv[2]
99+candidates = sorted(out_dir.glob(pattern), key=lambda p: p.stat().st_mtime)
100+if not candidates:
101+ raise SystemExit(f"no artifact matching {pattern} in {out_dir}")
102+print(candidates[-1])
103+PY
104+}
105+ 
106+maybe_clean_outputs() {
107+ local step="$1"
124 if [[ "${ENABLE_CLEAN}" == "TRUE" ]]; then108 if [[ "${ENABLE_CLEAN}" == "TRUE" ]]; then
125- echo "[1/5] Cleaning AscendC build outputs..."109+ echo "[${step}] Cleaning AscendC build outputs..."
126 (cd "${ASCENDC_ROOT}" && OIL_GAS_OPS_OUTPUT_DIR="${OUTPUT_DIR}" python3 build.py clean)110 (cd "${ASCENDC_ROOT}" && OIL_GAS_OPS_OUTPUT_DIR="${OUTPUT_DIR}" python3 build.py clean)
127 else111 else
128- echo "[1/5] Skip clean (use --clean to clean ascendc/_work and output/ first)."112+ echo "[${step}] Skip clean (use --clean to clean ascendc/_work and output/ first)."
129 fi113 fi
114+}
130 115 
131- echo "[2/5] Building AscendC OPP package (bash build.sh --pkg)..."116+build_opp_run() {
117+ local step="${1:-}"
118+ if [[ -n "${step}" ]]; then
119+ echo "[${step}] Building AscendC OPP package..."
120+ fi
132 local ops_item121 local ops_item
133 local ops_array=()122 local ops_array=()
134 while IFS= read -r ops_item; do123 while IFS= read -r ops_item; do
@@ -140,43 +129,34 @@ run_full_pipeline() {
140 build_args+=(-j "${JOBS}")129 build_args+=(-j "${JOBS}")
141 fi130 fi
142 (cd "${ASCENDC_ROOT}" && OIL_GAS_OPS_OUTPUT_DIR="${OUTPUT_DIR}" bash build.sh "${build_args[@]}")131 (cd "${ASCENDC_ROOT}" && OIL_GAS_OPS_OUTPUT_DIR="${OUTPUT_DIR}" bash build.sh "${build_args[@]}")
132+ RUN_PKG="$(latest_artifact "${OUTPUT_DIR}/${SOC_VERSION}" "*_opp_*.run")"
133+}
143 134 
144- local run_pkg135+install_opp_run() {
145- run_pkg="$(136+ local step="${1:-}"
146- python3 - <<'PY' "${OUTPUT_DIR}/${SOC_VERSION}"137+ if [[ -n "${step}" ]]; then
147-from pathlib import Path138+ echo "[${step}] Installing OPP package: ${RUN_PKG}"
148-import sys139+ fi
140+ bash "${RUN_PKG}" --quiet
141+}
149 142 
150-out_dir = Path(sys.argv[1])143+build_python_wheel() {
151-candidates = sorted(out_dir.glob("*_opp_*.run"), key=lambda p: p.stat().st_mtime)144+ local step="${1:-}"
152-if not candidates:145+ if [[ -n "${step}" ]]; then
153- raise SystemExit(f"no OPP run package found in {out_dir}")146+ echo "[${step}] Building oil_gas_ops_prospect wheel..."
154-print(candidates[-1])147+ fi
155-PY
156- )"
157- 
158- echo "[3/5] Installing OPP package: ${run_pkg}"
159- bash "${run_pkg}" --quiet
160- 
161- echo "[4/5] Building and installing oil_gas_ops_prospect wheel..."
162 rm -rf "${SCRIPT_DIR}/build" "${SCRIPT_DIR}"/*.egg-info148 rm -rf "${SCRIPT_DIR}/build" "${SCRIPT_DIR}"/*.egg-info
163 mkdir -p "${WHEEL_DIR}"149 mkdir -p "${WHEEL_DIR}"
164 (cd "${SCRIPT_DIR}" && python3 setup.py bdist_wheel --dist-dir "${WHEEL_DIR}")150 (cd "${SCRIPT_DIR}" && python3 setup.py bdist_wheel --dist-dir "${WHEEL_DIR}")
151+ WHEEL_PATH="$(latest_artifact "${WHEEL_DIR}" "oil_gas_ops_prospect-*.whl")"
152+ rm -rf "${SCRIPT_DIR}/build" "${SCRIPT_DIR}"/*.egg-info
153+}
165 154 
166- local wheel_path155+install_python_wheel() {
167- wheel_path="$(156+ local step="${1:-}"
168- python3 - <<'PY' "${WHEEL_DIR}"157+ if [[ -n "${step}" ]]; then
169-from pathlib import Path158+ echo "[${step}] Installing wheel: ${WHEEL_PATH}"
170-import sys159+ fi
171- 
172-wheel_dir = Path(sys.argv[1])
173-candidates = sorted(wheel_dir.glob("oil_gas_ops_prospect-*.whl"), key=lambda p: p.stat().st_mtime)
174-if not candidates:
175- raise SystemExit(f"no oil_gas_ops_prospect wheel found in {wheel_dir}")
176-print(candidates[-1])
177-PY
178- )"
179- 
180 python3 -m pip uninstall -y oil-gas-ops-prospect >/dev/null 2>&1 || true160 python3 -m pip uninstall -y oil-gas-ops-prospect >/dev/null 2>&1 || true
181 python3 - <<'PY'161 python3 - <<'PY'
182from pathlib import Path162from pathlib import Path
@@ -203,29 +183,73 @@ for root in {Path(p) for p in roots if p}:
203 for path in root.glob("complex_mul_npu*.so"):183 for path in root.glob("complex_mul_npu*.so"):
204 path.unlink(missing_ok=True)184 path.unlink(missing_ok=True)
205PY185PY
206- python3 -m pip install --force-reinstall --no-deps "${wheel_path}"186+ python3 -m pip install --force-reinstall --no-deps "${WHEEL_PATH}"
207- rm -rf "${SCRIPT_DIR}/build" "${SCRIPT_DIR}"/*.egg-info187+}
208 188 
209- if [[ "${ENABLE_SKIP_SMOKE}" != "TRUE" ]]; then189+run_whl_smoke() {
210- echo "[5/5] Running whl smoke tests..."190+ local step="$1"
211- # OIL_GAS_OPS_WHL_TEST=1 让 tests/conftest.py 不要把仓内源码挂到导入名下。191+ if [[ "${ENABLE_SKIP_SMOKE}" == "TRUE" ]]; then
212- OIL_GAS_OPS_WHL_TEST=1 \192+ echo "[${step}] Skip tests (--skip-tests)."
213- python3 -m pytest "${SCRIPT_DIR}/tests/test_oil_gas_ops_prospect_whl.py" -q193+ return 0
214- else
215- echo "[5/5] Skip tests (--skip-tests)."
216 fi194 fi
195+ echo "[${step}] Running whl smoke tests..."
196+ # OIL_GAS_OPS_WHL_TEST=1 让 tests/conftest.py 不要把仓内源码挂到导入名下。
197+ OIL_GAS_OPS_WHL_TEST=1 \
198+ python3 -m pytest "${SCRIPT_DIR}/tests/test_oil_gas_ops_prospect_whl.py" -q
199+}
217 200 
201+print_build_summary() {
202+ local installed="$1"
203+ echo
204+ echo "[oil_gas_ops_prospect] Done."
205+ echo "OPP package : ${RUN_PKG}"
206+ if [[ -n "${WHEEL_PATH}" ]]; then
207+ echo "Wheel : ${WHEEL_PATH}"
208+ fi
209+ if [[ "${installed}" != "TRUE" ]]; then
210+ echo "Install : skipped (packages left in output/)"
211+ return 0
212+ fi
218 cat <<EOF213 cat <<EOF
219 214 
220-[oil_gas_ops_prospect] Done.
221-OPP package : ${run_pkg}
222-Wheel : ${wheel_path}
223- 
224For later shells, keep this runtime library path:215For later shells, keep this runtime library path:
225export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:\$LD_LIBRARY_PATH216export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:\$LD_LIBRARY_PATH
226EOF217EOF
227}218}
228 219 
220+# --pkg: 只编 OPP .run。默认: 编 .run 与 whl,不安装。--install: 编译安装一体并冒烟。
221+run_build_pipeline() {
222+ echo "[oil_gas_ops_prospect] SOC_VERSION=${SOC_VERSION}"
223+ echo "[oil_gas_ops_prospect] OPS_LIST=${OPS_LIST}"
224+ 
225+ RUN_PKG=""
226+ WHEEL_PATH=""
227+ 
228+ if [[ "${ENABLE_PACKAGE}" == "TRUE" ]]; then
229+ maybe_clean_outputs "1/2"
230+ build_opp_run "2/2"
231+ print_build_summary FALSE
232+ return 0
233+ fi
234+ 
235+ if [[ "${ENABLE_INSTALL}" != "TRUE" ]]; then
236+ maybe_clean_outputs "1/3"
237+ build_opp_run "2/3"
238+ build_python_wheel "3/3"
239+ print_build_summary FALSE
240+ return 0
241+ fi
242+ 
243+ maybe_clean_outputs "1/5"
244+ build_opp_run "2/5"
245+ install_opp_run "3/5"
246+ echo "[4/5] Building and installing oil_gas_ops_prospect wheel..."
247+ build_python_wheel
248+ install_python_wheel
249+ run_whl_smoke "5/5"
250+ print_build_summary TRUE
251+}
252+ 
229main() {253main() {
230 checkopts "$@"254 checkopts "$@"
231 255 
@@ -239,11 +263,20 @@ main() {
239 echo "[ut] cleaning ophost UT build dir..."263 echo "[ut] cleaning ophost UT build dir..."
240 rm -rf "${ASCENDC_ROOT}/_work/ophost_ut"264 rm -rf "${ASCENDC_ROOT}/_work/ophost_ut"
241 fi265 fi
242- run_device_free_ut266+ run_unified_tests
243 return 0267 return 0
244 fi268 fi
245 269 
246- run_full_pipeline270+ if [[ -n "${OPENSOURCE_REPORT}" ]]; then
271+ local json_path="${SCRIPT_DIR}/test_result/test_results.json"
272+ echo "[oil_gas_ops_prospect] generating opensource report from ${json_path}"
273+ python3 "${SCRIPT_DIR}/scripts/opensource_report.py" \
274+ --from-json="${json_path}" \
275+ --out="${OPENSOURCE_REPORT}"
276+ return 0
277+ fi
278+ 
279+ run_build_pipeline
247}280}
248 281 
249main "$@"282main "$@"
@@ -6,11 +6,11 @@
6 6 
7| 项 | 要求 |7| 项 | 要求 |
8|----|------|8|----|------|
9-| CANN | 9.0.0(Toolkit 包 + 对应 SoC 的 Ops 包) |9+| CANN | 9.0.0 官方版本(Toolkit 包 + 对应 SoC 的 Ops 包) |
10| SoC | Ascend 910B(`Ascend910B1`) |10| SoC | Ascend 910B(`Ascend910B1`) |
11| Python | >= 3.10 |11| Python | >= 3.10 |
12-| Python 依赖 | `torch`、`torch-npu`、`triton`(Triton-Ascend),版本需与 CANN 配套 |12+| Python 依赖 | `torch`、`torch-npu`,以及 **Triton-Ascend 3.2.1**(不要用 PyPI 社区 `triton` 代替) |
13-| NPU 设备 | **精度与性能验证必须有真实 NPU**;编译与部分单元测试无卡即可 |13+| NPU 设备 | 真实 910B;精度、性能、显存必须在真机上验证通过 |
14 14 
15配置 CANN 环境变量:15配置 CANN 环境变量:
16 16 
@@ -18,6 +18,8 @@
18source /usr/local/Ascend/cann-9.0.0/set_env.sh18source /usr/local/Ascend/cann-9.0.0/set_env.sh
19```19```
20 20 
21+Triton-Ascend 安装步骤见[环境部署 · 安装 Triton-Ascend](zh/context/quick_install.md#6-安装-triton-ascend)。
22+ 
21## 2. 获取源码23## 2. 获取源码
22 24 
23```bash25```bash
@@ -25,15 +27,17 @@ git clone https://gitcode.com/cann/oil-gas-ops-prospect.git
25cd oil-gas-ops-prospect27cd oil-gas-ops-prospect
26```28```
27 29 
28-## 3. 一键构建并安装30+## 3. 编译与安装
29 31 
30-在仓库根目录执行:32+仓库根目录执行。默认只编译、不安装:
31 33 
32```bash34```bash
33-bash build.sh35+bash build.sh --pkg # 只出 OPP .run
36+bash build.sh # 出 .run 与 Python whl 到 output/,不安装
37+bash build.sh --install # 编译并安装,然后跑 whl 冒烟
34```38```
35 39 
36-该脚本依次完成:40+`--install` 依次完成:
37 41 
381. 调用 `ascendc/build.sh --pkg` 编译 AscendC OPP `.run` 包(默认 `Ascend910B1`,默认全部算子)。421. 调用 `ascendc/build.sh --pkg` 编译 AscendC OPP `.run` 包(默认 `Ascend910B1`,默认全部算子)。
392. 安装生成的 `oil_gas_ops_prospect_opp_*.run` 到当前 CANN 环境。432. 安装生成的 `oil_gas_ops_prospect_opp_*.run` 到当前 CANN 环境。
@@ -87,39 +91,19 @@ y = real_irfft(xr, xi, n=256, dim=-1) # 重建实信号
87 91 
88更多示例见 [`examples/`](../examples/) 与[算子调用](zh/invocation/quick_op_invocation.md)。92更多示例见 [`examples/`](../examples/) 与[算子调用](zh/invocation/quick_op_invocation.md)。
89 93 
90-## 5. 无卡验证94+## 5. 测试
91 95 
92-开发容器(如 `ascendc-dev`)没有 NPU 设备时,下列验证**不需要卡**:96+分类见 [`tests/README.md`](../tests/README.md)。在已安装 OPP / whl 的 910B 上执行:
93 97 
94```bash98```bash
95-# 全量无卡单元测试(--ut 是 -u 的别名)
96bash build.sh -u99bash build.sh -u
97-# 只跑 ophost / 只编译不执行:
98-# bash build.sh -u --ophost
99-# bash build.sh -u --noexec
100```100```
101 101 
102-等价的分步命令:102+算子在 NPU 上的精度、性能、显存必须在真实 910B 上验证通过。
103- 
104-```bash
105-# 1. ophost 侧 tiling 单元测试(googletest)
106-cmake -S . -B build && cmake --build build -j && ctest --test-dir build --output-on-failure
107-# 或:bash build.sh -u --ophost
108-# 或走 AscendC 侧入口:bash ascendc/build.sh -u --ophost
109- 
110-# 2. Triton 算子 CPU 回退精度
111-python3 -m pytest tests/triton/real_fft -q
112- 
113-# 3. Triton kernel 全分支 AOT 编译(TTIR → Linalg → bishengir → npubin)
114-python3 triton-ascend/tools/aot_compile_check.py
115-```
116- 
117-> **编译通过不等于功能正确。** 算子在 NPU 上的精度、性能、显存必须在真实 910B 上验证。
118 103 
119## 下一步104## 下一步
120 105 
121- [算子列表](zh/op_list.md):当前支持哪些算子、约束是什么106- [算子列表](zh/op_list.md):当前支持哪些算子、约束是什么
122- [接口列表](zh/api_list.md):完整接口签名与环境变量107- [接口列表](zh/api_list.md):完整接口签名与环境变量
123- [算子开发指南](zh/develop/operator_development_guide.md):新增算子步骤108- [算子开发指南](zh/develop/operator_development_guide.md):新增算子步骤
124-- [算子开发规格](zh/develop/operator-dev-spec.md):约束、交付件与验收
125- [贡献指南](../CONTRIBUTING.md):提交 Issue / PR109- [贡献指南](../CONTRIBUTING.md):提交 Issue / PR
@@ -7,7 +7,7 @@ oil-gas-ops-prospect 的完整文档索引。
7| 文档 | 说明 |7| 文档 | 说明 |
8|------|------|8|------|------|
9| [快速入门](QUICKSTART.md) | 从环境准备到跑通第一个算子的最短路径 |9| [快速入门](QUICKSTART.md) | 从环境准备到跑通第一个算子的最短路径 |
10-| [环境部署](zh/context/quick_install.md) | CANN Toolkit / Ops 包安装与环境变量配置 |10+| [环境部署](zh/context/quick_install.md) | CANN Toolkit / Ops 包、Triton-Ascend 安装与环境变量配置 |
11| [目录结构](zh/context/dir_structure.md) | 仓库各目录职责说明 |11| [目录结构](zh/context/dir_structure.md) | 仓库各目录职责说明 |
12 12 
13## 使用13## 使用
@@ -22,11 +22,10 @@ oil-gas-ops-prospect 的完整文档索引。
22 22 
23| 文档 | 说明 |23| 文档 | 说明 |
24|------|------|24|------|------|
25-| [编译参数说明](zh/context/build.md) | `build.sh` / `scripts/build_options.sh` / `ascendc/build.sh` 全部选项 |25+| [编译参数说明](zh/context/build.md) | `build.sh` 的编译构建与测试报告 |
26-| [算子开发指南](zh/develop/operator_development_guide.md) | 新增 AscendC / Triton 算子的步骤 |26+| [算子开发指南](zh/develop/operator_development_guide.md) | 新增 AscendC / Triton 算子的步骤、约定与验收 |
27-| [算子开发规格](zh/develop/operator-dev-spec.md) | 必须遵守的约束、交付件与验收 |
28| [精度验收报告模板](zh/develop/precision_acceptance_template.md) | PR 中附带的精度报告格式 |27| [精度验收报告模板](zh/develop/precision_acceptance_template.md) | PR 中附带的精度报告格式 |
29-| [调试与性能分析](zh/debug/op_debug_prof.md) | 常见问题定位与 profiling |28+| [调试与性能分析](zh/debug/op_debug_prof.md) | Host / Kernel 调试与 msProf 性能采集 |
30 29 
31## 社区30## 社区
32 31 
@@ -1,3 +0,0 @@
1-# 本文已迁移
2- 
3-算子开发与提交规范已迁到 [`zh/develop/operator-dev-spec.md`](zh/develop/operator-dev-spec.md),并按当前仓库布局刷新。请以新文档为准。
@@ -279,7 +279,7 @@ Triton 算子没有 aclnn 入口。
279|----------|------|------|------|279|----------|------|------|------|
280| `OIL_GAS_OPS_RUN_PKG` | 路径 | 自动探测 | `tests/ascendc/*/run_test.sh` 指定要安装的 OPP `.run` 包 |280| `OIL_GAS_OPS_RUN_PKG` | 路径 | 自动探测 | `tests/ascendc/*/run_test.sh` 指定要安装的 OPP `.run` 包 |
281| `OIL_GAS_OPS_WHL_REQUIRE_NPU_EXT` | `0` / `1` | `0` | 置 `1` 时 whl 冒烟测试强制要求已编译的 NPU 扩展存在,否则该用例跳过 |281| `OIL_GAS_OPS_WHL_REQUIRE_NPU_EXT` | `0` / `1` | `0` | 置 `1` 时 whl 冒烟测试强制要求已编译的 NPU 扩展存在,否则该用例跳过 |
282-| `OIL_GAS_OPS_WHL_TEST` | `0` / `1` | `0` | 由 `tests/run_whl_tests.sh` / `build.sh` 置 `1`,声明本次确实在验证已安装的 whl;此时不把仓内源码挂到 `oil_gas_ops_prospect` 导入名下,若解析到仓内路径则硬失败 |282+| `OIL_GAS_OPS_WHL_TEST` | `0` / `1` | `0` | 由 `tests/run_whl_tests.sh` / `bash build.sh --install` 置 `1`,声明本次确实在验证已安装的 whl;此时不把仓内源码挂到 `oil_gas_ops_prospect` 导入名下,若解析到仓内路径则硬失败 |
283| `OIL_GAS_OPS_RUN_LATENT3D_FFT_TEST` | `0` / `1` | `0` | 置 `1` 时运行 `(2,41,41,41,2048)` 大 shape 的 NPU FFT 重型用例 |283| `OIL_GAS_OPS_RUN_LATENT3D_FFT_TEST` | `0` / `1` | `0` | 置 `1` 时运行 `(2,41,41,41,2048)` 大 shape 的 NPU FFT 重型用例 |
284| `OIL_GAS_OPS_RUN_128_FFT_TEST` | `0` / `1` | `0` | 置 `1` 时运行 `tests/triton/real_fft` 的 128³ NPU 反向与显存用例 |284| `OIL_GAS_OPS_RUN_128_FFT_TEST` | `0` / `1` | `0` | 置 `1` 时运行 `tests/triton/real_fft` 的 128³ NPU 反向与显存用例 |
285 285 
@@ -1,75 +1,154 @@
1# 编译参数说明1# 编译参数说明
2 2 
3-本仓构建入口职责如下:3+## 简介
4 4 
5-| 入口 | 职责 |5+`build.sh` 位于仓库根目录,把源码编译成可安装的 OPP `.run` 与 Python whl,或按分类执行测试并写出报告。选项解析在 [`scripts/build_options.sh`](../../../scripts/build_options.sh)。
6-|------|------|
7-| `build.sh`(仓库根目录) | 一键完成 OPP 编译 → 安装 → whl 构建 → whl 安装 → 冒烟测试;`-u` 模式跑无卡单元测试 |
8-| `scripts/build_options.sh` | 被根目录 `build.sh` source:选项解析、互斥校验与分节帮助(对齐 [ops-math](https://gitcode.com/cann/ops-math/blob/master/scripts/build_options.sh)) |
9-| `ascendc/build.sh` | AscendC 侧分步入口:只打包 OPP `.run`,或只跑 ophost UT |
10-| `CMakeLists.txt`(顶层 CMake) | 无卡 C++ 目标(ophost tiling UT)的标准 CMake 构建 |
11 6 
12-## `build.sh`7+脚本按用途分成两条互斥路径,一次只走其中一条:
8+ 
9+| 路径 | 怎么进入 | 做什么 |
10+|------|----------|--------|
11+| [编译构建](#编译构建) | 不加 `-u` | 编包到 `output/`;加 `--install` 才安装 |
12+| [测试报告](#测试报告) | `-u` | 不编包、不安装;只跑测并写报告 |
13+ 
14+`--pkg`、`--install` 不能与 `-u` 同时使用,彼此也不能组合。本仓没有 ops-math 的 `--opapi` / `--opgraph` / `--opkernel`。
15+ 
16+## 使用方法
17+ 
18+1. 配置环境变量
19+ 
20+ 环境部署见[环境部署](quick_install.md)。默认路径安装、以 root 用户为例:
21+ 
22+ ```bash
23+ source /usr/local/Ascend/cann-9.0.0/set_env.sh
24+ ```
25+ 
26+ CANN 不在默认路径时:
27+ 
28+ ```bash
29+ export ASCEND_HOME_PATH=/path/to/cann
30+ source ${ASCEND_HOME_PATH}/set_env.sh
31+ ```
32+ 
33+2. 查看帮助
34+ 
35+ ```bash
36+ bash build.sh --help # 总览
37+ bash build.sh --pkg -h # 编译构建分节
38+ bash build.sh --install -h # 同上
39+ bash build.sh -u -h # 测试报告分节
40+ bash build.sh --clean -h # 清理分节
41+ ```
42+ 
43+ `--soc=SOC` / `--soc SOC`、`--ops=a,b` / `--ops a,b`、`-j16` / `-j 16` 均支持。全量参数含义见下方两节,按实际场景选用。
44+ 
45+---
46+ 
47+## 编译构建
48+ 
49+不加 `-u` 时走构建路径。根目录 `build.sh` 是主入口;`ascendc/build.sh` 与顶层 CMake 见节末[其它入口](#其它入口)。
50+ 
51+### 构建命令格式
52+ 
53+以编译并安装为例。`--soc`、`--ops`、`-j` 在该场景为可选项,合适的编译线程数可加快编译速度。
13 54 
14```bash55```bash
15-bash build.sh [options]56+bash build.sh --install [--soc=${soc_version}] [--ops=${op_list}] [-j${n}]
16-bash build.sh -h # 总览
17-bash build.sh -u -h # Test Options
18-bash build.sh --clean -h # Clean Options
19```57```
20 58 
21-选项解析在 [`scripts/build_options.sh`](../../../scripts/build_options.sh)。`--soc=SOC` / `--soc SOC`、`--ops=a,b` / `--ops a,b`、`-j16` / `-j 16` 均支持。59+### 三种入口
22 60 
23-### 默认(打包安装)61+三种入口互斥,一次只选一种。
24 62 
25-不加 `-u` 时走全流程:OPP → 安装 → whl → 冒烟。63+**单独编译 OPP `.run` 包**
26- 
27-| 选项 | 默认 | 说明 |
28-|------|------|------|
29-| `--soc=<SOC>` | `Ascend910B1` | msopgen SOC 型号 |
30-| `--ops=op1,op2,...` | `all` | 逗号或空格分隔的 AscendC 算子名 |
31-| `-j[n]`, `--jobs <N>` | CPU 核数 | 编译并行度,透传给 `ascendc/build.py` |
32-| `--clean` | 关 | 构建前清理 `ascendc/_work` 与 `output/` |
33-| `--skip-tests` | 关 | 安装但跳过 whl 冒烟测试 |
34-| `-h, --help` | — | 打印帮助;与 `-u` / `--clean` 组合打印对应分节 |
35- 
36-`--pkg` 不能与 `-u` 同时使用(不加 `-u` 本来就是打包路径)。
37 64 
38```bash65```bash
39-bash build.sh # 全流程66+bash build.sh --pkg
40-bash build.sh --clean # 先清理再全流程67+```
68+ 
69+只调用 `ascendc/build.sh --pkg` 生成 OPP `.run`,写入 `output/<SOC>/`。不编 Python whl,不安装到 CANN。
70+ 
71+**只编译不安装(默认)**
72+ 
73+```bash
74+bash build.sh
75+```
76+ 
77+编译 OPP `.run` 与 Python whl,产物写入 `output/`,不执行 `.run` 安装器、也不 `pip install`。编 whl 仍需链接 CANN 中已有的 `libcust_opapi.so`(与直接 `python3 setup.py bdist_wheel` 相同);全新环境请先 `bash build.sh --install`,或先安装已有的 `.run`。
78+ 
79+**编译并安装**
80+ 
81+```bash
82+bash build.sh --install
83+```
84+ 
85+编译 OPP 与 whl,安装到 CANN / Python,并跑 whl 冒烟测试。安装完成后,后续 shell 需要保留自定义 opapi 动态库路径:
86+ 
87+```bash
88+export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:$LD_LIBRARY_PATH
89+```
90+ 
91+### 参数说明
92+ 
93+`bash build.sh --help` 可查看全部构建参数。
94+ 
95+| 参数名 | 可选/必选 | 参数说明 |
96+|--------|-----------|----------|
97+| `--pkg` | 可选 | 只编译 OPP `.run`,不编 whl、不安装。不可与 `--install`、`-u` 同时使用。 |
98+| `--install` | 可选 | 编译 `.run` 与 whl,安装后跑冒烟。不可与 `--pkg`、`-u` 同时使用。 |
99+| (默认,无上述开关) | — | 编译 `.run` 与 whl 到 `output/`,不安装。 |
100+| `--soc=${soc}` | 可选 | 指定 NPU 型号,默认 `Ascend910B1`。每次编译只支持 1 个型号。 |
101+| `--ops=${op_list}` | 可选 | 指定待编译的 AscendC 算子,逗号或空格分隔,默认全部。 |
102+| `-j${n}`, `--jobs <N>` | 可选 | 编译线程数,默认 CPU 核数(如 `-j16`)。 |
103+| `--clean` | 可选 | 构建前清理 `ascendc/_work` 与 `output/`,然后继续执行所选入口。 |
104+| `--skip-tests` | 可选 | 仅配合 `--install`:安装但跳过 whl 冒烟测试。 |
105+| `-h`, `--help` | 可选 | 打印帮助;与 `--pkg` / `--install` / `--clean` 组合打印对应分节。 |
106+ 
107+### 示例
108+ 
109+```bash
110+bash build.sh --pkg # 只出 OPP .run
111+bash build.sh # 出 .run 与 whl,不安装
112+bash build.sh --install # 编译并安装
113+bash build.sh --clean # 先清理再编译(默认仍不安装)
41bash build.sh --ops=complex_mul,complex_mul_grad # 只构建指定算子114bash build.sh --ops=complex_mul,complex_mul_grad # 只构建指定算子
42-bash build.sh --skip-tests # 跳过冒烟测试115+bash build.sh --install --skip-tests # 安装后跳过冒烟测试
43-bash build.sh --soc=Ascend910B1 -j16 # 指定 SOC 与并行度116+bash build.sh --soc=Ascend910B1 -j16 --install # 指定 SOC 与并行度并安装
44```117```
45 118 
46-### Test Options(`-u`)119+### 产物位置
47 120 
48-对齐 ops-math 的 `Test Options`:`-u` 跑全部无卡 UT;加上 `--ophost` / `--triton` / `--aot` 则只跑对应子集。`--ut` 是 `-u` 的别名。121+所有编译产物统一落到仓库根目录 `output/`:
49 122 
50-| 选项 | 说明 |123+```text
51-|------|------|124+output/
52-| `-u`, `--ut` | 编译并运行全部无卡单元测试(不打包不安装) |125+├── <SOC>/oil_gas_ops_prospect_opp_<os>_<arch>.run # OPP 算子包
53-| `--noexec` | 只编译 ophost UT,不执行;Triton / AOT 同步跳过 |126+└── wheels/oil_gas_ops_prospect-*.whl # Python whl
54-| `--cov` | 打开覆盖率:ophost 用 `--coverage`;Triton 在已装 `pytest-cov` 时加 `--cov` |127+```
55-| `--gtest_filter=pattern` | 只跑匹配的 ophost gtest(需配合 `-u` 或 `-u --ophost`) |128+ 
56-| `--ophost -u` | 只跑 ophost tiling UT(googletest) |129+设置 `OIL_GAS_OPS_OUTPUT_DIR` 可改变该位置,对 `build.sh`、`ascendc/build.sh --pkg`、`tests/run_whl_tests.sh` 均生效。
57-| `--triton -u` | 只跑 Triton CPU 回退 pytest |130+ 
58-| `--aot -u` | 只跑 Triton AOT 编译检查 |131+| 交付物 | 名称 | 规范 |
132+|--------|------|------|
133+| Python 分发名 | `oil-gas-ops-prospect` | PEP 508 |
134+| import 包名 | `oil_gas_ops_prospect` | PEP 8 |
135+| whl | `oil_gas_ops_prospect-<ver>-<py>-<abi>-<plat>.whl` | PEP 427(`-`/`.` 归一为 `_`) |
136+| OPP 算子包 | `oil_gas_ops_prospect_opp_<os>_<arch>.run` | CANN `<pkg>_opp_<os>_<arch>.run` |
137+| 环境变量前缀 | `OIL_GAS_OPS_` | — |
138+ 
139+CANN vendor 目录仍为约定的 `opp/vendors/customize/`。
140+ 
141+不需要重新编译 AscendC 算子、只验证已安装的 whl 时:
59 142 
60```bash143```bash
61-bash build.sh -u # 全部无卡 UT144+bash tests/run_whl_tests.sh
62-bash build.sh -u --noexec # 只编译 ophost UT
63-bash build.sh -u --ophost # 只跑 ophost
64-bash build.sh -u --ophost --gtest_filter=ComplexMul*
65-bash build.sh -u --triton --aot
66-bash build.sh --ophost --triton --aot -u # 与上一行等价(-u 可放后面)
67-bash build.sh --ut # 兼容旧写法,等同 -u
68```145```
69 146 
70-`--ophost` / `--triton` / `--aot` 必须搭配 `-u`。本仓没有 ops-math 的 `--opapi` / `--opgraph` / `--opkernel` UT。147+### 其它入口
71 148 
72-## `ascendc/build.sh`149+日常使用根目录 `build.sh` 即可。下面两个入口给需要分步调试或只编 ophost tiling 单元测试的场景。
150+ 
151+**`ascendc/build.sh`**:AscendC 侧分步入口,只打包 OPP `.run`,或只跑 ophost UT。
73 152 
74```bash153```bash
75cd ascendc154cd ascendc
@@ -79,35 +158,21 @@ bash build.sh clean # 清理 ascendc/_work 与 output/
79bash build.sh list-ops # 列出 manifest 中的算子158bash build.sh list-ops # 列出 manifest 中的算子
80```159```
81 160 
82-打包相关选项(透传给 `build.py`):161+| 参数名 | 可选/必选 | 参数说明 |
83- 162+|--------|-----------|----------|
84-| 选项 | 默认 | 说明 |163+| `-v`, `--soc-version <SOC>` | 可选 | msopgen SOC 型号,默认 `Ascend910B1`,支持逗号分隔的多型号 |
85-|------|------|------|164+| `--ops <ops...>` | 可选 | 算子目录名,可多个或用逗号分隔,默认全部 |
86-| `-v, --soc-version <SOC>` | `Ascend910B1` | msopgen SOC 型号,支持逗号分隔的多型号 |165+| `-j`, `--jobs <N>` | 可选 | 编译并行度,默认 CPU 核数 |
87-| `--ops <ops...>` | 全部 | 算子目录名,可多个或用逗号分隔 |166+| `-i`, `--install-path <path>` | 可选 | CANN 安装路径,默认自动探测 |
88-| `-j, --jobs <N>` | CPU 核数 | 编译并行度 |167+| `--vendor-name <name>` | 可选 | OPP 安装目录 `opp/vendors/<name>/`,默认 `customize` |
89-| `-i, --install-path <path>` | 自动探测 | CANN 安装路径 |168+| `--package-name <name>` | 可选 | 替换 `.run` 名中的 `custom` 前缀,默认 `oil_gas_ops_prospect` |
90-| `--vendor-name <name>` | `customize` | OPP 安装目录 `opp/vendors/<name>/` |
91-| `--package-name <name>` | `oil_gas_ops_prospect` | 替换 `.run` 名中的 `custom` 前缀 |
92- 
93-示例:
94 169 
95```bash170```bash
96bash build.sh --pkg -v Ascend910B1 --ops complex_mul complex_mul_grad -j 16171bash build.sh --pkg -v Ascend910B1 --ops complex_mul complex_mul_grad -j 16
97-bash build.sh --pkg --vendor-name customize --package-name oil_gas_ops_prospect172+python3 build.py -v Ascend910B1 --ops complex_mul complex_mul_grad # 等价 Python 入口
98```173```
99 174 
100-等价的 Python 入口:175+**顶层 CMake**:只覆盖 ophost 侧 tiling 单元测试。AscendC 算子本身不由顶层 CMake 编译,必须经 msopgen 生成工程后再由 `bash build.sh --pkg` 或 `bash ascendc/build.sh --pkg` 编译。
101- 
102-```bash
103-python3 build.py -v Ascend910B1 --ops complex_mul complex_mul_grad
104-python3 build.py clean
105-python3 build.py list-ops
106-```
107- 
108-## 顶层 CMake
109- 
110-顶层 CMake 只覆盖**无卡即可构建**的 C++ 目标,即 ophost 侧 tiling 单元测试:
111 176 
112```bash177```bash
113cmake -S . -B build -DCMAKE_BUILD_TYPE=Release178cmake -S . -B build -DCMAKE_BUILD_TYPE=Release
@@ -122,46 +187,92 @@ ctest --test-dir build --output-on-failure
122| `OIL_GAS_GTEST_VERSION` | `1.14.0` | googletest 版本 |187| `OIL_GAS_GTEST_VERSION` | `1.14.0` | googletest 版本 |
123| `OIL_GAS_GTEST_URL` | gitcode 镜像地址 | googletest 源码包地址,仅在本地找不到 gtest 时使用 |188| `OIL_GAS_GTEST_URL` | gitcode 镜像地址 | googletest 源码包地址,仅在本地找不到 gtest 时使用 |
124 189 
125-**AscendC 算子本身不由顶层 CMake 编译**:CANN 自定义算子必须由 msopgen 依据当前 CANN 与 SOC 生成配套工程后,再交给 opbuild / bisheng 编译,入口是 `bash build.sh` 与 `bash ascendc/build.sh --pkg`。190+版本与依赖定义集中在 [`version.cmake`](../../../version.cmake):`OIL_GAS_OPS_PROSPECT_VERSION`(需与 `pyproject.toml` / `setup.py` 一致)、`OIL_GAS_OPS_PROSPECT_CANN_VERSION`、`OIL_GAS_OPS_PROSPECT_SOC_VERSIONS`、以及构建期 / 运行期依赖变量。
126 191 
127-版本与依赖定义集中在 [`version.cmake`](../../../version.cmake):192+---
128 193 
129-| 变量 | 说明 |194+## 测试报告
130-|------|------|
131-| `OIL_GAS_OPS_PROSPECT_VERSION` | 仓库版本,需与 `pyproject.toml` / `setup.py` 一致 |
132-| `OIL_GAS_OPS_PROSPECT_CANN_VERSION` | 已验证的 CANN 版本 |
133-| `OIL_GAS_OPS_PROSPECT_SOC_VERSIONS` | 已验证的 SoC 型号 |
134-| `OIL_GAS_OPS_PROSPECT_BUILD_DEPENDENCIES` | 构建期依赖 |
135-| `OIL_GAS_OPS_PROSPECT_RUN_DEPENDENCIES` | 运行期依赖 |
136 195 
137-## 产物位置196+`-u` 进入测试模式:不编包、不安装。分类与用例路径以 [`tests/README.md`](../../../tests/README.md) 为准。本节能直接写出两类产物:
138 197 
139-所有编译产物统一落到仓库根目录 `output/`:198+| 产物 | 何时生成 | 路径 |
199+|------|----------|------|
200+| 执行纪要 | 每次 `-u` | `test_result/test_report.md`(另存带时间戳的副本)与 `test_result/test_results.json` |
201+| 开源测试报告 | 加 `--opensource-report` | `test_result/oil-gas-ops-prospect仓开源测试报告.md`,并覆盖 `docs/zh` 同名文件 |
140 202 
141-```text203+### 测试命令格式
142-output/
143-├── <SOC>/oil_gas_ops_prospect_opp_<os>_<arch>.run # OPP 算子包
144-└── wheels/oil_gas_ops_prospect-*.whl # Python whl
145-```
146 204 
147-设置 `OIL_GAS_OPS_OUTPUT_DIR` 可改变该位置,对 `build.sh`、`ascendc/build.sh --pkg`、`tests/run_whl_tests.sh` 均生效。205+不跟子集时跑全部(含 FFT 大形态与开源门禁)。`--ops`、`--report`、`--opensource-report` 为可选项。
148- 
149-## 只构建并验证 whl
150- 
151-不需要重新编译 AscendC 算子时:
152 206 
153```bash207```bash
154-bash tests/run_whl_tests.sh208+bash build.sh -u [--ut|--st|--precision|--perf|...] [--ops=${op_list}] [--opensource-report]
155```209```
156 210 
157-## 包命名规范211+无 NPU 时,全量 `-u` 的上板项 **SKIP**,不失败。显式指定 `--ut` / `--st` / `--precision` / `--perf` / `--examples` 且无 NPU 时 **FAIL**。
158 212 
159-| 交付物 | 名称 | 规范 |213+`--ut` / `--st` / `--precision` / `--perf` 等子集开关本身也会进入测试模式,因此 `bash build.sh --ut` 与 `bash build.sh -u --ut` 等价。CI `JOB_llt` 走 cmake/ctest,不调用 `build.sh -u`。
160-|--------|------|------|
161-| Python 分发名 | `oil-gas-ops-prospect` | PEP 508 |
162-| import 包名 | `oil_gas_ops_prospect` | PEP 8 |
163-| whl | `oil_gas_ops_prospect-<ver>-<py>-<abi>-<plat>.whl` | PEP 427(`-`/`.` 归一为 `_`) |
164-| OPP 算子包 | `oil_gas_ops_prospect_opp_<os>_<arch>.run` | CANN `<pkg>_opp_<os>_<arch>.run` |
165-| 环境变量前缀 | `OIL_GAS_OPS_` | — |
166 214 
167-CANN vendor 目录仍为约定的 `opp/vendors/customize/`。215+### 参数说明
216+ 
217+`bash build.sh -u -h` 可查看测试分节帮助。
218+ 
219+| 参数名 | 可选/必选 | 对应类别 | 参数说明 |
220+|--------|-----------|----------|----------|
221+| `-u` | 进入测试模式 | 全部 | 不跟子集时跑下表全部 |
222+| `--all` | 可选 | 全部 | 显式全量,与裸 `-u` 相同 |
223+| `--ut` | 可选 | **UT** | UT-host(切分)+ UT-device(上板小 launch);需要 NPU |
224+| `--ophost` | 可选 | UT-host | 只跑 tiling googletest |
225+| `--st` | 可选 | **ST** | 多 shape C++ aclnn + golden;需要 NPU |
226+| `--precision` | 可选 | **精度** | NPU Python 对拍 |
227+| `--perf` | 可选 | **性能** | NPU Python 计时 |
228+| `--examples` | 可选 | UT-device | 只跑 C++ 小 launch;需要 NPU |
229+| `--triton` | 可选 | Triton CPU | CPU 回退 pytest |
230+| `--aot` | 可选 | Triton AOT | kernel 可编译检查 |
231+| `--ops=${op_list}` | 可选 | 过滤 | 作用于 UT-device / ST / 精度 / 性能 |
232+| `--report=<path>` | 可选 | — | 执行纪要 Markdown 路径 |
233+| `--opensource-report[=path]` | 可选 | — | 按本次 `-u` 的 JSON 生成开源测试报告;单独使用时只读已有 JSON、不再跑测 |
234+| `--noexec` | 可选 | UT-host | 只编译,不执行 |
235+| `--cov` | 可选 | UT-host | 打开 `--coverage`;Triton 在已装 pytest-cov 时加 `--cov` |
236+| `--gtest_filter=pattern` | 可选 | UT-host | 过滤 gtest |
237+| `--clean` | 可选 | — | 跑测前清理 `ascendc/_work` 与 `output/` |
238+ 
239+`--noexec` / `--cov` / `--gtest_filter` / `--report` 只能与测试模式一起用。`--opensource-report` 不能与 `--pkg` / `--install` 组合。
240+ 
241+### 示例
242+ 
243+```bash
244+bash build.sh -u # 全部
245+bash build.sh -u --ut # UT(切分 + 上板)
246+bash build.sh -u --st --precision --perf
247+bash build.sh -u --ophost # 只切分
248+bash build.sh -u --perf --ops=complex_mul
249+bash build.sh -u --report=test_result/my.md # 自定义执行纪要路径
250+bash build.sh -u --opensource-report # 跑测并生成开源测试报告
251+bash build.sh --opensource-report # 基于上次 -u 的 JSON 只出报告
252+```
253+ 
254+### 执行纪要
255+ 
256+每次 `bash build.sh -u` 跑完都会落盘:
257+ 
258+```text
259+test_result/
260+├── test_report_<时间戳>.md # 本次执行纪要
261+├── test_report.md # 上一份纪要的稳定副本
262+└── test_results.json # 开源测试报告的数据源
263+```
264+ 
265+`--report=<path>` 只改 Markdown 纪要路径,不改变 JSON 位置。
266+ 
267+### 开源测试报告
268+ 
269+社区模板开源测试报告**不**随裸 `-u` 自动生成,需要显式加 `--opensource-report`。
270+ 
271+| 命令 | 行为 |
272+|------|------|
273+| `bash build.sh -u --opensource-report` | 跑齐测试(含文档 / 出包入口 / 安装痕迹 / 接口冒烟,以及 FFT 大形态),再按本次 JSON **生成**报告 |
274+| `bash build.sh --opensource-report` | 只读已有 `test_result/test_results.json`,不再跑测 |
275+ 
276+默认写出 `test_result/oil-gas-ops-prospect仓开源测试报告.md`,并同步覆盖 `docs/zh` 同名文件。有 FAIL 时结论为不建议发布,不会回填旧数字。
277+ 
278+分类表、套件目录与新增算子应挂哪些测试,见 [`tests/README.md`](../../../tests/README.md)。
@@ -13,8 +13,8 @@ oil-gas-ops-prospect/
13├── .gitignore # Git 忽略规则13├── .gitignore # Git 忽略规则
14├── .clang-format # C++ 代码格式14├── .clang-format # C++ 代码格式
15├── .pre-commit-config.yaml # 提交前检查钩子15├── .pre-commit-config.yaml # 提交前检查钩子
16-├── build.sh # 顶层构建入口(OPP + whl + 冒烟;-u 跑无卡 UT)16+├── build.sh # 顶层构建入口(默认只编包;--install 安装;-u 跑测试)
17-├── CMakeLists.txt # 顶层 CMake(无卡 C++ 目标:ophost UT)17+├── CMakeLists.txt # 顶层 CMake(ophost tiling UT)
18├── version.cmake # 版本与依赖定义18├── version.cmake # 版本与依赖定义
19├── requirements.txt # Python 依赖声明19├── requirements.txt # Python 依赖声明
20├── pyproject.toml # Python 打包元数据20├── pyproject.toml # Python 打包元数据
@@ -24,7 +24,7 @@ oil-gas-ops-prospect/
24│ └── workflows/ # GitCode Action 门禁流水线24│ └── workflows/ # GitCode Action 门禁流水线
25│ ├── PR-pipeline_oil-gas-ops-prospect.yml # 门禁总编排25│ ├── PR-pipeline_oil-gas-ops-prospect.yml # 门禁总编排
26│ ├── compile_action.yml # 编译(x86 / aarch64)26│ ├── compile_action.yml # 编译(x86 / aarch64)
27-│ ├── llt_action.yml # 无卡 LLT(ophost UT + Triton 精度 + AOT 编译)27+│ ├── llt_action.yml # LLT(ophost UT + Triton 精度 + AOT 编译)
28│ └── staticcheck_action.yml # 静态检查28│ └── staticcheck_action.yml # 静态检查
29│29│
30├── cmake/ # CMake 模块30├── cmake/ # CMake 模块
@@ -39,7 +39,7 @@ oil-gas-ops-prospect/
39│ ├── operators/39│ ├── operators/
40│ │ ├── manifest.tsv # 算子注册表(tab 分隔)40│ │ ├── manifest.tsv # 算子注册表(tab 分隔)
41│ │ ├── _common/ # 多算子共用的 host/kernel 头41│ │ ├── _common/ # 多算子共用的 host/kernel 头
42-│ │ ├── complex_mul/ # ComplexMul42+│ │ ├── complex_mul/ # ComplexMul(仅 op_host / op_kernel)
43│ │ │ ├── op_host/ # OpDef / InferShape / Tiling43│ │ │ ├── op_host/ # OpDef / InferShape / Tiling
44│ │ │ └── op_kernel/ # kernel 实现44│ │ │ └── op_kernel/ # kernel 实现
45│ │ ├── complex_mul_grad/ # ComplexMulGrad45│ │ ├── complex_mul_grad/ # ComplexMulGrad
@@ -60,22 +60,31 @@ oil-gas-ops-prospect/
60│ │ ├── triton_fft_real.py # Triton kernel 与 autograd 封装60│ │ ├── triton_fft_real.py # Triton kernel 与 autograd 封装
61│ │ └── README.md61│ │ └── README.md
62│ └── tools/62│ └── tools/
63-│ └── aot_compile_check.py # 无卡 AOT 编译检查63+│ └── aot_compile_check.py # AOT 编译检查
64│64│
65-├── tests/ # ★ 全部测试用例(算子目录内不放用例)65+├── tests/ # ★ 全部测试(operators/ 不再放 tests)
66+│ ├── suites.json # build.sh -u 套件清单
67+│ ├── _common/aclnn/ # ST 共用脚手架
68+│ ├── ut/op_host/<op>/ # tiling googletest
69+│ ├── st/aclnn<OpType>/ # 多 case 两段式 aclnn
66│ ├── conftest.py # pytest 收集配置70│ ├── conftest.py # pytest 收集配置
67│ ├── test_oil_gas_ops_prospect_whl.py # 基于已安装 whl 的冒烟用例71│ ├── test_oil_gas_ops_prospect_whl.py # 基于已安装 whl 的冒烟用例
68│ ├── run_whl_tests.sh # 构建 whl → 安装 → 跑冒烟72│ ├── run_whl_tests.sh # 构建 whl → 安装 → 跑冒烟
69│ ├── ascendc/73│ ├── ascendc/
70-│ │ ├── ophost/ # tiling UT(googletest,无卡可跑)74+│ │ ├── ophost/ # 汇总编译 ophost_ut
71│ │ ├── complex_mul/ # NPU 精度/性能对比测试75│ │ ├── complex_mul/ # NPU 精度/性能对比测试
72│ │ └── fused_bias_softmax/ # 融合注意力 NPU 精度/性能对比76│ │ └── fused_bias_softmax/ # 融合注意力 NPU 精度/性能对比
73│ └── triton/77│ └── triton/
74-│ └── real_fft/ # rFFT / iRFFT 精度(无卡回退 CPU)78+│ └── real_fft/ # rFFT / iRFFT 精度(含 CPU 回退路径)
75│79│
76-├── examples/ # 示例代码80+├── examples/ # Python 示例 + C++ 两段式样例
81+│ ├── _common/ # aclnn 样例脚手架
82+│ ├── aclnn/<op>/ # test_aclnn_*.cpp
83+│ └── *.py
77├── scripts/ # 辅助脚本84├── scripts/ # 辅助脚本
78│ ├── build_options.sh # 根目录 build.sh 的选项解析与分节帮助85│ ├── build_options.sh # 根目录 build.sh 的选项解析与分节帮助
86+│ ├── run_tests.py # 统一编排 UT/ST/精度/性能并写 Markdown 报告
87+│ ├── opensource_report.py # 按 -u JSON 生成开源测试报告
79│ ├── check_license_header.sh # 许可证头检查88│ ├── check_license_header.sh # 许可证头检查
80│ └── local_pkg.py # 裸 checkout 下解析 oil_gas_ops_prospect 导入名89│ └── local_pkg.py # 裸 checkout 下解析 oil_gas_ops_prospect 导入名
81├── docs/ # 文档90├── docs/ # 文档
@@ -85,9 +94,8 @@ oil-gas-ops-prospect/
85│ ├── op_list.md # 算子列表94│ ├── op_list.md # 算子列表
86│ ├── api_list.md # 接口列表95│ ├── api_list.md # 接口列表
87│ ├── context/ # 背景与环境96│ ├── context/ # 背景与环境
88-│ ├── develop/ # 开发指南与规格97+│ ├── develop/ # 开发指南
89-│ │ ├── operator_development_guide.md # 新增算子步骤98+│ │ ├── operator_development_guide.md # 新增算子步骤、约定与验收
90-│ │ ├── operator-dev-spec.md # 约束、交付件与验收
91│ │ └── precision_acceptance_template.md99│ │ └── precision_acceptance_template.md
92│ ├── debug/ # 调试与性能100│ ├── debug/ # 调试与性能
93│ └── invocation/ # 调用示例101│ └── invocation/ # 调用示例
@@ -4,11 +4,11 @@
4 4 
5当前仓库已验证通过的 CANN 版本如下:5当前仓库已验证通过的 CANN 版本如下:
6 6 
7-| CANN 版本 | 发布时间 | 分支 |7+| CANN 版本 | 发布时间 |
8-|-----------|----------|------|8+|-----------|----------|
9-| [CANN 9.0.0](https://www.hiascend.com/developer/download/community/result?module=cann&cann=9.0.0) | 2026/04/30 | master |9+| [9.0.0 官方版本](https://www.hiascend.com/developer/download/community/result?module=cann&cann=9.0.0) | 2026/04/30 |
10 10 
11-本仓库需要安装 **Toolkit 包**和对应 SoC 的 **Ops 包**。请根据实际 CPU 架构和产品型号,从上述链接目录中选择对应的 `.run` 安装包。11+本仓库需要安装 **Toolkit 包**、对应 SoC 的 **Ops 包**,以及与 CANN 配套的 **[Triton-Ascend](#6-安装-triton-ascend)**。请根据实际 CPU 架构和产品型号,从上述链接目录中选择对应的 `.run` 安装包。
12 12 
13安装包文件名格式:13安装包文件名格式:
14 14 
@@ -69,14 +69,33 @@ source ${install_path}/cann/set_env.sh
69pip install -r requirements.txt69pip install -r requirements.txt
70```70```
71 71 
72-> `torch-npu` 与 Triton-Ascend 需与本地 CANN 版本严格配套,通常随 CANN 配套发布渠道安装,不建议直接从 PyPI 拉取任意版本。72+> `torch-npu` 需与本地 CANN 版本严格配套,通常随 CANN 配套发布渠道安装,不建议直接从 PyPI 拉取任意版本。
73 73 
74-## 6. 安装本项目74+## 6. 安装 Triton-Ascend
75 75 
76-见[快速入门](../../QUICKSTART.md)第 3 节,或直接:76+本仓的 `real_rfft` / `real_irfft` 依赖 **Triton-Ascend** 提供的 `triton.backends.ascend` 后端。PyPI 上的社区 `triton` 包没有昇腾后端,不能代替下面的安装。本仓已验证 **CANN 9.0.0 官方版本** 配套 **`triton-ascend==3.2.1`**。
77+ 
78+上游说明见 [Triton-Ascend](https://gitcode.com/Ascend/triton-ascend)([安装指南](https://triton-ascend.readthedocs.io/zh-cn/latest/installation_guide.html))。安装命令:
77 79 
78```bash80```bash
79-bash build.sh81+pip uninstall -y triton triton-ascend
82+pip install triton-ascend==3.2.1 --extra-index-url=https://mirrors.huaweicloud.com/ascend/repos/pypi
83+```
84+ 
85+安装后确认包名是 `triton-ascend`,且版本为 `3.2.1`:
86+ 
87+```bash
88+python3 -c "import importlib.metadata as m; print(m.version('triton-ascend'))"
89+```
90+ 
91+若之后再 `pip install` 其它依赖把社区 `triton` 装回来,按上面两条命令重装即可。本仓不下载、不分发 Triton-Ascend。
92+ 
93+## 7. 安装本项目
94+ 
95+见[快速入门 · 编译与安装](../../QUICKSTART.md#3-编译与安装),或直接:
96+ 
97+```bash
98+bash build.sh --install
80```99```
81 100 
82安装完成后,后续 shell 中需要保留自定义 opapi 动态库路径:101安装完成后,后续 shell 中需要保留自定义 opapi 动态库路径:
@@ -85,12 +104,8 @@ bash build.sh
85export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:$LD_LIBRARY_PATH104export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:$LD_LIBRARY_PATH
86```105```
87 106 
88-## 7. 容器环境说明
89- 
90-在开发容器(如 `ascendc-dev`)中,仓库通常挂载在 `/root/projects/oil-gas-ops-prospect`。容器内一般没有 NPU 设备,此时只能做编译与无卡验证,见[快速入门](../../QUICKSTART.md)第 5 节。
91- 
92如果 CANN 安装路径不是默认的 `/usr/local/Ascend/cann-9.0.0`,通过环境变量指定:107如果 CANN 安装路径不是默认的 `/usr/local/Ascend/cann-9.0.0`,通过环境变量指定:
93 108 
94```bash109```bash
95-ASCEND_HOME_PATH=/path/to/cann bash build.sh110+ASCEND_HOME_PATH=/path/to/cann bash build.sh --install
96```111```
@@ -1,44 +1,140 @@
1# 调试与性能分析1# 调试与性能分析
2 2 
3-## 常见问题定位3+本仓 AscendC 算子跑在 **AI Core**(`ComplexMul` / `ComplexMulGrad` 为 AIV,`FusedSoftmaxGrad` 的 bf16 路径为 MIX)。文档说明本仓能用来定位问题、核对精度和采集性能的入口。
4+ 
5+精度、性能与显存必须在真实 **910B(NPU)** 上验证通过:
6+ 
7+```bash
8+bash build.sh -u
9+```
10+ 
11+---
12+ 
13+## 调试定位(AI Core 算子)
14+ 
15+算子执行失败或精度异常时,可分别看 Host 日志和 Kernel 中间结果。
16+ 
17+### 1、Host 侧日志
18+ 
19+程序结束后,Host 调试日志默认在:
20+ 
21+```text
22+$HOME/ascend/log/debug/plog/plog-pid_*.log
23+```
24+ 
25+路径与分类见 CANN 9.0 [查看日志](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900/maintenref/logreference/logreference_0002.html)。
26+ 
27+开启 [`ASCEND_SLOG_PRINT_TO_STDOUT`](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900/maintenref/envvar/envref_07_0121.html) 可将 Host 应用类日志打到标准输出(`1` 开启,`0` 关闭)。开启后日志不再写入 log 文件:
28+ 
29+```bash
30+export ASCEND_SLOG_PRINT_TO_STDOUT=1
31+```
32+ 
33+需要更详细级别时设置 [`ASCEND_GLOBAL_LOG_LEVEL`](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900/maintenref/envvar/envref_07_0001.html)(仅对调试日志生效:`0` DEBUG、`1` INFO、`2` WARNING、`3` ERROR):
34+ 
35+```bash
36+export ASCEND_GLOBAL_LOG_LEVEL=1
37+```
38+ 
39+完整环境变量列表见 [环境变量参考](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900/maintenref/envvar/envref_07_0001.html)。
40+ 
41+Python 调用路径建议打开同步执行,让报错栈指向真实出错位置:
42+ 
43+```python
44+import torch_npu
45+torch_npu.npu.set_sync_debug_mode(1)
46+```
47+ 
48+aclnn 调用失败时,用 [`aclGetRecentErrMsg`](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900/API/runtimeapi/aclcppdevg_03_0028.html) 取最近一次 ACL 错误描述(取完即清空,应在失败后立刻调用)。本仓样例 `examples/aclnn/complex_mul/test_aclnn_complex_mul.cpp` 的 `LaunchOp` 目前只打印返回码,排查时可改成:
49+ 
50+```cpp
51+auto ret = aclnnComplexMulGetWorkspaceSize(
52+ t.inT[0], t.inT[1], t.inT[2], t.inT[3], t.wrT, t.wiT, &workspaceSize, &executor);
53+if (ret != ACL_SUCCESS) {
54+ const char* msg = aclGetRecentErrMsg();
55+ printf("aclnnComplexMulGetWorkspaceSize failed. ERROR: %d\n%s\n",
56+ ret, msg == nullptr ? "" : msg);
57+ return ret;
58+}
59+```
60+ 
61+头文件已由 `examples/_common/aclnn_example.h` 引入 `acl/acl.h`。
62+ 
63+### 2、Kernel 调试
64+ 
65+`AscendC::printf` / `DumpTensor` 会拖慢 Kernel,只在调测阶段打开。接口说明见 [printf](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900/API/ascendcopapi/atlasascendc_api_07_0193.html) 与 [DumpTensor](https://www.hiascend.com/document/detail/zh/canncommercial/900/API/ascendcopapi/atlasascendc_api_07_0192.html)。
66+ 
67+本仓样例里 `aclInit(nullptr)`(见 `examples/_common/aclnn_example.h` 的 `Init`),默认不导出 Kernel dump。调测时在可执行文件目录放 `acl.json`,并把 `aclInit(nullptr)` 改成 `aclInit("acl.json")`:
68+ 
69+```json
70+{
71+ "dump": {
72+ "dump_kernel_data": "printf,tensor",
73+ "dump_path": "./kernel_dump"
74+ }
75+}
76+```
77+ 
78+落盘路径优先级:`ASCEND_DUMP_PATH` > `ASCEND_WORK_PATH` > 配置里的 `dump_path`。导出的 bin 不能直接当文本看,用 [show_kernel_debug_data](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900/programug/Ascendcopdevg/atlas_ascendc_10_0102.html) 解析;`printf` 内容在 `parser.log`,`DumpTensor` 在 `dump_data/`。
79+ 
80+- printf:打印标量。在 `ComplexMul` 读完 tiling 后(`op_kernel/complex_mul.cpp` 的 `complex_mul` 入口)可临时加上:
81+ 
82+```c++
83+ComplexMulTilingData tilingData;
84+OgOpsKernel::CopyTilingData(&tilingData, tiling);
85+AscendC::printf("totalElems=%ld tileFloats=%ld coreNum=%ld blockIdx=%d\n",
86+ tilingData.totalElems, tilingData.tileFloats, tilingData.coreNum,
87+ static_cast<int>(GetBlockIdx()));
88+```
89+ 
90+- DumpTensor:打印 UB 上的 LocalTensor。在 `ComplexMulKernel::ProcessTile` 写出 `wr` 之前(`op_kernel/complex_mul_impl.h`)可临时加上,第三参是元素个数,单核 dump 总量不要超过 1MB:
91+ 
92+```c++
93+out = outQue_.DeQue<float>();
94+const uint32_t dumpN = static_cast<uint32_t>(cur < 32 ? cur : 32);
95+DumpTensor(out, 0, dumpN);
96+CopyOutPad(wrGm_[gmOff], out, cur);
97+```
98+ 
99+卡死、GM / UB 越界等场景可用 [msDebug](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900beta2/devaids/optool/atlasopdev_16_0062.html) 单步调试。
100+ 
101+---
102+ 
103+## 本仓常见问题
4 104 
5### 自定义算子没有被调用105### 自定义算子没有被调用
6 106 
7-现象:设置了 `OIL_GAS_OPS_COMPLEX_MUL_IMPL=ascendc`,但性能与 PyTorch 实现一致。107+现象:设了 `OIL_GAS_OPS_COMPLEX_MUL_IMPL=ascendc`(或融合算子的 `OIL_GAS_OPS_FUSED_BIAS_SOFTMAX_IMPL=ascendc`),但结果/耗时仍像 PyTorch 参考实现。
8 108 
9-原因:`complex_mul` 在 AscendC 调用失败时**默认静默回退**到 PyTorch 实现。109+原因:AscendC 调用失败时默认静默回退。禁止回退:
10- 
11-定位:
12 110 
13```bash111```bash
14export OIL_GAS_OPS_REQUIRE_ASCENDC_COMPLEX_MUL=1112export OIL_GAS_OPS_REQUIRE_ASCENDC_COMPLEX_MUL=1
113+# 融合注意力算子:
114+export OIL_GAS_OPS_REQUIRE_ASCENDC_FUSED_BIAS_SOFTMAX=1
15```115```
16 116 
17-置 1 后调用失败会直接抛异常,暴露真实原因。常见根因:117+环境变量全集见[接口列表](../api_list.md#环境变量)。
18 118 
19| 根因 | 排查 |119| 根因 | 排查 |
20|------|------|120|------|------|
21-| pybind 扩展 `complex_mul_npu` 未编译或未安装 | `python3 -c "import complex_mul_npu"` |121+| pybind 扩展未编译或未安装 | `python3 -c "import complex_mul_npu"`(融合算子为 `fused_bias_softmax_npu`) |
22| OPP 包未安装 | `ls ${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib` |122| OPP 包未安装 | `ls ${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib` |
23-| 动态库路径缺失 | `echo $LD_LIBRARY_PATH`,确认含 `opp/vendors/customize/op_api/lib` |123+| 动态库路径缺失 | `echo $LD_LIBRARY_PATH`,需含 `opp/vendors/customize/op_api/lib` |
24| 张量不在 NPU 上 | `x.device.type == "npu"` |124| 张量不在 NPU 上 | `x.device.type == "npu"` |
25-| dtype 不是 float32 或几个输入 shape 不一致 | 检查入参 |125+| dtype / shape 不满足算子约束 | 见[算子列表](../op_list.md) |
26- 
27-补齐动态库路径:
28 126 
29```bash127```bash
30-export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:$LD_LIBRARY_PATH128+export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:${LD_LIBRARY_PATH}
31```129```
32 130 
33### 找不到 aclnn 接口131### 找不到 aclnn 接口
34 132 
35-自定义算子的 aclnn 接口名随 OpType 变化。本仓 `ComplexMul` 对应 `aclnnComplexMul`。改了 OpType 就要同步改 pybind 绑定。133+自定义算子的 aclnn 名随 OpType 变化。本仓 `ComplexMul` 对应 `aclnnComplexMul`。改了 OpType 必须同步改 `ascendc/pybind/csrc/` 里的绑定。
36 134 
37### `.run` 包安装后不生效135### `.run` 包安装后不生效
38 136 
39-OPP 包安装到 `opp/vendors/<vendor_name>/`,默认 `customize`。若 CANN 侧配置了多个 vendor,需确认 `opp/vendors/config.ini` 中的优先级包含 `customize`。137+OPP 安装到 `opp/vendors/<vendor_name>/`,默认 `customize`。多个 vendor 时确认 `opp/vendors/config.ini` 含 `customize`。
40- 
41-重装:
42 138 
43```bash139```bash
44bash output/Ascend910B1/oil_gas_ops_prospect_opp_*.run --quiet140bash output/Ascend910B1/oil_gas_ops_prospect_opp_*.run --quiet
@@ -46,73 +142,114 @@ bash output/Ascend910B1/oil_gas_ops_prospect_opp_*.run --quiet
46 142 
47### 编译报错:缺少 CANN 内部框架头143### 编译报错:缺少 CANN 内部框架头
48 144 
49-从 CANN 原仓移植的算子常依赖 `op_host/tiling_base.h`、`err/ops_err.h`、`log/log.h` 等内部头,msopgen 自定义算子工程的 opbuild 环境下不可用。处理方式见[算子开发指南](../develop/operator_development_guide.md)第 8 节。145+从 CANN 原仓移植的 host 代码常依赖 `op_host/tiling_base.h`、`err/ops_err.h`、`log/log.h` 等,msopgen 自定义工程里没有这些头。处理方式见[算子开发指南 · 移植 CANN 框架算子](../develop/operator_development_guide.md#10-移植-cann-框架算子的注意事项)。
50 146 
51### Triton 算子在 CPU 上跑但没报错147### Triton 算子在 CPU 上跑但没报错
52 148 
53-`real_rfft` / `real_irfft` 在非 NPU 环境(或 Triton 不可用)会**自动回退到 `torch.fft`**。这是设计行为,用于无卡精度自检。确认是否走了 NPU kernel,检查张量 device 与 Triton 是否可用。149+`real_rfft` / `real_irfft` 在非 NPU 或 Triton 不可用时回退到 `torch.fft`。要确认走了 NPU kernel,检查张量 device 以及 `triton-ascend` 是否已安装。
54 150 
55-## 无卡验证手段151+---
56 152 
57-容器(如 `ascendc-dev`)没有 NPU 时,下列手段不需要卡:153+## 性能采集
154+ 
155+时延或显存异常时,用 `msprof op` 看 Kernel 耗时、核数和流水占比。下面命令已在真实 910B、CANN 9.0.0 上用本仓 `ComplexMul` 样例跑通。其它算子把可执行文件换成 `examples/aclnn/<op>/output/test_aclnn_<op>`,仿真 `--kernel-name` 换成该算子 OpType。
156+ 
157+- **上板采集**:真实 NPU 上拿 Kernel 耗时、Block 数、流水占比。
158+- **流水图仿真**:看各核指令流水;比上板更细。
159+ 
160+该样例是 `{2,2}` 小 tensor,上板 `Task Duration` 只有数微秒,用来验证采集链路,不能当业务性能数字。代表 shape 的 `ms/call` 仍走本节末尾的 Python `--perf`。
161+ 
162+命令与结果文件含义见 [获取性能数据](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/900/programug/Ascendcopdevg/atlas_ascendc_best_practices_10_0008.html)。
163+ 
164+### 方式 1:上板性能采集
165+ 
166+先安装 customize vendor,再编样例。`run.sh` 会清掉 `build/`、`output/`,重新 cmake/make,并**立刻跑一遍**功能检查:
58 167 
59```bash168```bash
60-# 全量无卡单元测试169+source /usr/local/Ascend/cann-9.0.0/set_env.sh
61-bash build.sh -u170+export ASCEND_HOME_PATH=/usr/local/Ascend/cann-9.0.0
171+export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:${LD_LIBRARY_PATH}
172+bash examples/aclnn/complex_mul/run.sh
62```173```
63 174 
64-分项说明:175+再在输出目录做性能采集(这是第二次 launch;默认预热 5 次):
65- 
66-| 手段 | 命令 | 覆盖什么 |
67-|------|------|----------|
68-| ophost tiling UT | `bash ascendc/build.sh -u --ophost` | 切分算法:UB 预算约束、对齐与上限、退化输入、核间切分覆盖性不变量 |
69-| 顶层 CMake + ctest | `cmake -S . -B build && cmake --build build -j && ctest --test-dir build --output-on-failure` | 同上,走标准 CMake 流程 |
70-| Triton CPU 回退精度 | `python3 -m pytest tests/triton/real_fft -q` | Triton 算子语义(对 `torch.fft` 对拍) |
71-| Triton AOT 编译 | `python3 triton-ascend/tools/aot_compile_check.py` | 全部 kernel 分支可编译(TTIR → Linalg → bishengir → npubin) |
72-| OPP 打包 | `bash ascendc/build.sh --pkg` | 算子能被 opbuild / bisheng 编译并打包 |
73-| 许可证头 | `bash scripts/check_license_header.sh` | 合规 |
74- 
75-> **编译通过不等于功能正确。** 精度、性能、显存必须在真实 910B 上验证。
76- 
77-## NPU 上的精度与性能验证
78 176 
79```bash177```bash
80-cd tests/ascendc/complex_mul178+cd examples/aclnn/complex_mul/output
81-bash run_test.sh179+msprof op ./test_aclnn_complex_mul
82- 
83-# 指定 shape
84-bash run_test.sh --shape 1 1 96 160 160
85- 
86-# 指定 OPP 包
87-OIL_GAS_OPS_RUN_PKG=/path/to/oil_gas_ops_prospect_opp_openEuler_aarch64.run bash run_test.sh
88```180```
89 181 
90-`run_test.sh` 分两步:安装/更新 OPP 包 → 编译 pybind 扩展并跑对比脚本。也可手动分步:182+结果目录为同级 `OPPROF_<时间戳>_*`。本仓实测摘要:
91- 
92-```bash
93-bash ../../../output/Ascend910B1/oil_gas_ops_prospect_opp_openEuler_aarch64.run --quiet
94-python3 test_complex_mul_npu.py --build-ext # 编译扩展并测试
95-python3 test_complex_mul_npu.py # 复用已编译扩展
96-```
97- 
98-输出形如:
99 183 
100```text184```text
101-[fwd] shape=... rel_err wr=... wi=...185+Op Name: ComplexMul_923973e39f803f5bc6ddb804acabf320_0
102-[bwd] rel_err gur=... gui=... gvr=... gvi=...186+Op Type: vector
103-[perf] fwd torch=...ms ascendc=...ms speedup=...x187+Task Duration(us): 4.640093
104-[perf] bwd torch=...ms ascendc=...ms speedup=...x188+Block Dim: 48
105-[autograd] gur=... gui=... gvr=... gvi=...189+Device Id: 0
106-PASS
107```190```
108 191 
109-退出码:`0` 精度达标;`1` 超阈值 / 无 NPU / 扩展缺失。192+`Task Duration` 是当前 Kernel 耗时,`Block Dim` 是执行核数。本仓实测结果目录(时间戳与随机串每次不同):
110 193 
111-## 性能调优194+```text
195+OPPROF_<时间戳>_<随机串>/
196+├── OpBasicInfo.csv # 核名、Op Type、总耗时、Block Dim、设备号
197+├── ArithmeticUtilization.csv # 各 block 的 cube/vector 算力占比(vector 看 aiv_time / aiv_vec_ratio)
198+├── PipeUtilization.csv # 各流水(VEC / MTE2 / MTE3 / SCALAR)耗时与占比
199+├── Memory.csv # GM ↔ UB 带宽、搬运量
200+├── MemoryL0.csv # L0A/L0B/L0C 带宽(纯 vector 核多为 NA)
201+├── MemoryUB.csv # UB 读写带宽
202+├── L2Cache.csv # L2 命中 / miss
203+├── ResourceConflictRatio.csv # bank / MTE 冲突与 wait 占比
204+└── dump/
205+ ├── op_basic_info.txt # 与 OpBasicInfo 对应的文本摘要
206+ └── pc_start_addr.txt # Kernel PC 起始地址
207+```
208+ 
209+`FusedSoftmaxGrad` 的 bf16 路径为 MIX,`Op Type` 可能不是纯 vector;该算子尚未用 `msprof op` 复测。
210+ 
211+Python 精度 / 性能对比(相对误差与 `ms/call`,含业务 shape):
212+ 
213+```bash
214+bash build.sh -u --precision --perf --ops=complex_mul
215+# 或(会尝试安装 output/ 下的 OPP .run,再编 pybind、跑对比)
216+cd tests/ascendc/complex_mul && bash run_test.sh
217+```
218+ 
219+### 方式 2:仿真流水图采集
220+ 
221+本仓默认 `SOC_VERSION=Ascend910B1`。仿真前把仿真库加入 `LD_LIBRARY_PATH`(按实际 CANN 安装路径与 SOC 目录修改)。本仓实测 `--kernel-name=ComplexMul` 即可命中上板同一核名,不必改成入口符号 `complex_mul`:
222+ 
223+```bash
224+export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/tools/simulator/Ascend910B1/lib:${LD_LIBRARY_PATH}
225+cd examples/aclnn/complex_mul/output
226+msprof op simulator --output=$PWD/pipeline_auto --kernel-name=ComplexMul ./test_aclnn_complex_mul
227+```
228+ 
229+结果在 `$PWD/pipeline_auto/OPPROF_*`(CANN 9.0.0 不落 `visualize_data.bin`)。本仓实测结构:
230+ 
231+```text
232+pipeline_auto/OPPROF_<时间戳>_<随机串>/
233+├── dump/
234+│ ├── pc_start_addr.txt # Kernel PC 起始地址
235+│ └── object_dump.txt # 仿真 dump 路径与核名
236+└── simulator/
237+ ├── trace.json # 全核 Chrome Tracing 事件(优先看这个)
238+ └── coreN.veccore0/ # 每个 vector 核一份;N 与 Block Dim 对应(本样例 0..47)
239+ ├── trace.json # 单核流水时间线
240+ ├── coreN.veccore0_instr_exe.csv # 指令级:流水、cycles、耗时
241+ └── coreN.veccore0_code_exe.csv # 源码行耗时;未编 -g 时只有表头
242+```
243+ 
244+用 Chrome `chrome://tracing` 或 MindStudio Insight 导入 `simulator/trace.json`。看某核堵在哪条指令时打开对应 `*_instr_exe.csv`。
245+ 
246+---
247+ 
248+## 本仓性能相关开关
112 249 
113### Triton FFT 分块250### Triton FFT 分块
114 251 
115-`OIL_GAS_OPS_FFT_BLOCK_M` 控制 DFT-matmul 的行分块大小(默认 512,受 L0C 中 fp32 累加器 ≤128KB 约束)。扫参:252+`OIL_GAS_OPS_FFT_BLOCK_M` 控制 DFT-matmul 的行分块(默认 512,受 910B L0C 中 fp32 累加器 ≤128KB 约束)。实现会按容量收缩,溢出时自动减半重试:
116 253 
117```bash254```bash
118for m in 128 256 512; do255for m in 128 256 512; do
@@ -120,28 +257,13 @@ for m in 128 256 512; do
120done257done
121```258```
122 259 
123-`verify_triton_fft_real_128_3d.py` 也支持 `--block-m` 参数覆盖该环境变量。260+`verify_triton_fft_real_128_3d.py` 的 `--block-m` 会覆盖该环境变量。
124 261 
125### AscendC 切分262### AscendC 切分
126 263 
127-切分算法在各算子的 `op_host/<op>_tiling_core.h`,改动后先用 ophost UT 验证不变量,再上板测性能。264+切分算法在各算子的 `ascendc/operators/<op>/op_host/<op>_tiling_core.h`。改完先跑 ophost UT,再在 910B 上板看性能:
128- 
129-### msprof 采集
130- 
131-CANN 提供的 `msprof` 可采集算子级性能数据,用于定位瓶颈(AIV 利用率、搬运/计算比、UB 命中)。采集需要真实 NPU 设备。
132- 
133-## 日志
134- 
135-CANN 运行日志(plog)默认在 `~/ascend/log/`。提高日志级别:
136 265 
137```bash266```bash
138-export ASCEND_GLOBAL_LOG_LEVEL=1 # 0-DEBUG 1-INFO 2-WARNING 3-ERROR267+bash build.sh -u --ophost
139-export ASCEND_SLOG_PRINT_TO_STDOUT=1268+bash build.sh -u --perf --ops=complex_mul
140-```
141- 
142-算子执行异常时建议同时开启同步执行,让报错栈指向真实出错位置:
143- 
144-```python
145-import torch_npu
146-torch_npu.npu.set_sync_debug_mode(1)
147```269```
@@ -1,460 +0,0 @@
1-# 算子开发与提交规范
2- 
3-> **适用范围**:本仓库新增或修改的 AscendC(`ascendc/operators/`)与 Triton-Ascend(`triton-ascend/`)算子。
4->
5-> **与[开发指南](operator_development_guide.md)的分工**:指南讲步骤;本文讲必须遵守的约束、交付件与验收。路径以 [`dir_structure.md`](../context/dir_structure.md) 为准。
6->
7-> **版本**:v2.0(按开源布局与当前实现刷新;不再沿用仓外 e3nn 的 `testplan/` / `teste2e/` / `delivery/` 目录约定)。
8- 
9----
10- 
11-## 目录
12- 
13-1. [总则](#1-总则)
14-2. [目录与交付件](#2-目录与交付件)
15-3. [AscendC 源码规范](#3-ascendc-源码规范)
16-4. [Python 封装规范](#4-python-封装规范)
17-5. [Triton-Ascend 规范](#5-triton-ascend-规范)
18-6. [构建与打包](#6-构建与打包)
19-7. [测试与精度](#7-测试与精度)
20-8. [PR 产物清单](#8-pr-产物清单)
21-9. [审核与迭代](#9-审核与迭代)
22-10. [分支与提交](#10-分支与提交)
23-11. [附录:检查清单](#11-附录检查清单)
24- 
25----
26- 
27-## 1. 总则
28- 
29-### 1.1 目标
30- 
31-每个算子从源码到真机运行形成可复现、可审核的闭环:编译通过不等于功能正确。
32- 
33-### 1.2 核心原则
34- 
35-- **闭环验证**:源码 → 编译 / 打包 → 安装 → 无卡 UT → 真机精度 / 性能 / 显存。
36-- **TDD**:先写 ophost UT 或 CPU 回退用例,再写实现。
37-- **双环境**:无卡覆盖切分与回退逻辑;NPU 覆盖精度、性能、内存安全。
38-- **测试不进算子目录**:用例一律放 `tests/`,`ascendc/operators/` 与 `triton-ascend/` 只放实现。
39- 
40-### 1.3 命名约定
41- 
42-| 侧 | 规则 | 示例 |
43-|----|------|------|
44-| C++ OpType | PascalCase | `ComplexMul`、`FusedBiasSoftmax` |
45-| aclnn | `aclnn` + OpType | `aclnnComplexMul` |
46-| kernel 入口 / `opFile.value` | 蛇形,与源文件名一致 | `complex_mul` |
47-| Python 公开函数 | 蛇形,**不加** `npu_` 前缀 | `complex_mul`、`fused_bias_softmax` |
48-| 反向 OPP 算子 | `Grad` 后缀 | `ComplexMulGrad`、`FusedSoftmaxGrad` |
49-| 环境变量 | `OIL_GAS_OPS_` 前缀 | `OIL_GAS_OPS_COMPLEX_MUL_IMPL` |
50- 
51-OpType 改名不影响 kernel 源码(入口由 `opFile.value` / `opInterface.value` 固定),但 aclnn 符号随 OpType 变化。
52- 
53-一对前反向可以共用一个 Python 模块(`fused_bias_softmax.py` 同时导出 `fused_bias_softmax` 与 `fused_softmax_grad`),但 **manifest 与 `operators/` 各占一行 / 一个目录**。
54- 
55-### 1.4 支持的 SoC
56- 
57-当前交付与真机验证面向 **Ascend 910B**(默认 `Ascend910B1`)。`ascendc/scripts/config.py` 的 `SUPPORTED_SOC_VERSIONS` 还可注册构建用 SOC;新增 SOC 必须同时:
58- 
59-1. 写入 `SUPPORTED_SOC_VERSIONS`;
60-2. 在 OpDef 的 `AICore().AddConfig(...)` 中增加对应配置;
61-3. 在该 SOC 上完成精度验收。
62- 
63----
64- 
65-## 2. 目录与交付件
66- 
67-### 2.1 仓库对照(当前布局)
68- 
69-```text
70-oil-gas-ops-prospect/
71-├── ascendc/
72-│ ├── operators/
73-│ │ ├── manifest.tsv
74-│ │ ├── _common/ # 多算子共用 host/kernel 头
75-│ │ └── <op_name>/op_host|op_kernel/
76-│ ├── pybind/ # 打包为 oil_gas_ops_prospect
77-│ │ ├── <op_family>.py
78-│ │ └── csrc/<op_family>/bindings.cpp
79-│ ├── pybind_extensions.py # 按 csrc/<op_family>/ 自动发现扩展
80-│ └── build.sh
81-├── triton-ascend/<op_family>/
82-├── tests/
83-│ ├── ascendc/ophost/ # googletest,无卡
84-│ ├── ascendc/<op_name>/ # NPU 精度/性能
85-│ └── triton/<op_family>/
86-├── setup.py # whl:package_dir 组装 + 自动带上 pybind 扩展
87-└── build.sh # 顶层:OPP + whl + 冒烟;-u 无卡 UT
88-```
89- 
90-产物落在仓库根 `output/`(可用 `OIL_GAS_OPS_OUTPUT_DIR` 覆盖),**不要**再写 `ascendc/dist/`、`python/`、`tests/`。
91- 
92-### 2.2 AscendC 单算子目录
93- 
94-以现有 `complex_mul` / `fused_bias_softmax` 为准:
95- 
96-```text
97-ascendc/operators/<op_name>/
98-├── op_host/
99-│ ├── <OpType>.cpp 或 <op_name>.cpp # OpDef + InferShape + TilingFunc 入口
100-│ ├── <op_name>_tiling.cpp # 【可选】gert 入口较长时拆出
101-│ ├── <op_name>_tiling_data_def.h # BEGIN_TILING_DATA_DEF
102-│ ├── <op_name>_tiling_core.h # 【必须】平台无关切分算法
103-│ └── <op_name>_tiling_impl.h # 对接 PlatformAscendC / TilingContext
104-└── op_kernel/
105- ├── <op_name>.cpp # kernel 入口
106- ├── <op_name>_impl.h # 【可选】实现
107- └── <op_name>_tiling.h # device 侧 plain struct 镜像
108-```
109- 
110-切分算法**必须**抽到 `*_tiling_core.h`:`TilingFunc` 依赖 `gert::TilingContext`,平台读取依赖 `PlatformAscendC`,CANN 二进制包没有 gert faker,UT 只能直接驱动 core。
111- 
112-### 2.3 一个目录对应一个 OpType
113- 
114-`manifest.tsv` 一行一个 OpType(tab 分隔:目录名、OpType、kernel 入口文件名):
115- 
116-```tsv
117-# 目录名 OpType(CANN) kernel入口文件名
118-complex_mul ComplexMul complex_mul.cpp
119-complex_mul_grad ComplexMulGrad complex_mul_grad.cpp
120-fused_bias_softmax FusedBiasSoftmax fused_bias_softmax.cpp
121-fused_softmax_grad FusedSoftmaxGrad fused_softmax_grad.cpp
122-```
123- 
124-前反向即使 Python 同模块,也要分两个算子目录。共享逻辑放 `_common/` 或各自的 `*_kernel.h`,入口 `.cpp` 分别 `#include`。
125- 
126----
127- 
128-## 3. AscendC 源码规范
129- 
130-### 3.1 OpDef(rt2.0)
131- 
132-使用 `ops::OpDef` + `OP_ADD`,不要用旧的 `REG_OP`。
133- 
134-```cpp
135-namespace ops {
136-class ComplexMul : public OpDef {
137-public:
138- explicit ComplexMul(const char *name) : OpDef(name)
139- {
140- DeclareRequiredFloatNdIo(*this, {"ur", "ui", "vr", "vi"}, {"wr", "wi"});
141- this->SetInferShape(ge::ComplexMulInferShape).SetInferDataType(ge::ComplexMulInferDataType);
142- this->AICore().SetTiling(optiling::ComplexMulTilingFunc);
143- AddAscend910bElementwiseConfig(*this, "complex_mul"); // opFile / opInterface
144- }
145-};
146-OP_ADD(ComplexMul);
147-}
148-```
149- 
150-**约定**:
151- 
152-- `opFile.value` / `opInterface.value` 必须等于 `op_kernel/` 下入口文件名(无扩展名)。
153-- 需要 aclnn 时设 `aclnnSupport.value = support_aclnn`(或依赖 msopgen 默认生成两段式入口,以该算子实测为准,见 `FusedSoftmaxGrad`)。
154-- 每个 OpType 单独一份 OpDef 源文件。
155-- 逐元素 float32/ND 算子优先复用 `_common/op_host/og_tiling_context.h` 的 `DeclareRequiredFloatNdIo` / `AddAscend910bElementwiseConfig`。
156- 
157-### 3.2 Tiling
158- 
159-Host 侧用 `BEGIN_TILING_DATA_DEF`(opbuild 需要);kernel 侧用 **plain struct** 镜像同一布局(`ccec` 不识别该宏)。
160- 
161-```cpp
162-// op_host/<op>_tiling_data_def.h
163-BEGIN_TILING_DATA_DEF(ComplexMulTilingDataDef)
164- TILING_DATA_FIELD_DEF(int64_t, totalElems);
165- TILING_DATA_FIELD_DEF(int64_t, tileFloats);
166- TILING_DATA_FIELD_DEF(int64_t, coreNum);
167-END_TILING_DATA_DEF;
168-REGISTER_TILING_DATA_CLASS(ComplexMul, ComplexMulTilingDataDef)
169-```
170- 
171-`TilingFunc` 负责:形状 / dtype 校验 → 调 `*_tiling_core.h` → `SetBlockDim` / `SetTilingKey` → `SaveToBuffer`。`GetCoreNumAiv()` 为 0 必须失败或给出明确兜底,并在 ophost UT 覆盖该分支。
172- 
173-### 3.3 Kernel 入口
174- 
175-```cpp
176-extern "C" __global__ __aicore__ void complex_mul(GM_ADDR ur, GM_ADDR ui, GM_ADDR vr, GM_ADDR vi,
177- GM_ADDR wr, GM_ADDR wi, GM_ADDR workspace, GM_ADDR tiling)
178-{
179- // 按 tiling 分核、分拍;AIV-only 算子在 kernel 侧声明 KERNEL_TYPE_AIV_ONLY
180-}
181-```
182- 
183-入口符号必须与 `opFile.value` 一致。
184- 
185-### 3.4 910B 向量 64 元素对齐
186- 
187-Ascend 910B 向量指令以 **64 个元素**为粒度。`count` 不是 64 倍数时,尾部可能被跳过或写成 0。
188- 
189-```cpp
190-uint32_t aligned = (count + 63u) & ~63u;
191-CompareScalar(mask, x, 0.0f, CMPMODE::GT, aligned); // 计算用 aligned
192-DataCopyPad(outGm_[offset], out, outParams); // 写出用原始 count,禁止越界写
193-```
194- 
195-`DataCopy` / `repeatTimes` 等其它对齐限制见 ascendc-api-best-practices;UB 按 32B 规划。
196- 
197-### 3.5 移植 CANN 框架算子
198- 
199-host tiling 若依赖 `tiling_base.h`、`ops_err.h`、`log/log.h` 等内部头,在 msopgen 工程里用本地 compat shim 替换宏,不要改切分算法。额外 include 目录加到 `ascendc/scripts/patch_op_host.py` 的 `_EXTRA_INCLUDE_SUBDIRS`。详见 [ascendc/README.md](../../../ascendc/README.md)。
200- 
201----
202- 
203-## 4. Python 封装规范
204- 
205-### 4.1 位置与发现
206- 
207-| 文件 | 作用 |
208-|------|------|
209-| `ascendc/pybind/<op_family>.py` | 公开 API,打包为 `oil_gas_ops_prospect.<op_family>` |
210-| `ascendc/pybind/csrc/<op_family>/bindings.cpp` | 直调 aclnn 两段式入口,编成 `<op_family>_npu` |
211- 
212-新建 `csrc/<op_family>/*.cpp` 后,`pybind_extensions.py` 会自动带上扩展,**不必**再手写 `setup.py` 的 `CppExtension`。Triton 子包仍需在 `setup.py` 的 `packages` / `package_dir` 登记。
213- 
214-### 4.2 标准模式
215- 
216-以现有 `complex_mul.py` 为准:
217- 
218-```python
219-@lru_cache(maxsize=1)
220-def _load_extension():
221- return importlib.import_module("complex_mul_npu")
222- 
223-def _require_custom() -> bool:
224- return os.environ.get("OIL_GAS_OPS_REQUIRE_ASCENDC_COMPLEX_MUL", "0") == "1"
225- 
226-def _impl() -> str:
227- return os.environ.get("OIL_GAS_OPS_COMPLEX_MUL_IMPL", "torch").strip().lower()
228- 
229-def torch_complex_mul(...): # 纯 PyTorch 参考,CPU / 对拍
230- ...
231- 
232-class _ComplexMulFn(torch.autograd.Function):
233- @staticmethod
234- def forward(ctx, ur, ui, vr, vi, use_ascendc):
235- ...
236- @staticmethod
237- def backward(ctx, gwr, gwi):
238- ...
239- 
240-def complex_mul(ur, ui, vr, vi):
241- use_ascendc = _impl() in ("ascendc", "auto") and ur.device.type == "npu"
242- return _ComplexMulFn.apply(ur, ui, vr, vi, use_ascendc)
243-```
244- 
245-### 4.3 必须遵守
246- 
247-| 规则 | 说明 |
248-|------|------|
249-| 函数式 API 走 `autograd.Function.apply` | 直接调 `ext.forward()` 会断梯度图 |
250-| 必须有独立 PyTorch 参考 | 默认 `IMPL=torch`;对拍与无卡测试依赖它 |
251-| 实现开关 | `OIL_GAS_OPS_<OP>_IMPL=torch\|ascendc\|auto`,默认 `torch` |
252-| 强制自定义算子 | `OIL_GAS_OPS_REQUIRE_ASCENDC_<OP>=1` 时禁止静默回退 |
253-| `@lru_cache` 缓存扩展 | 避免重复 import |
254-| 公开名蛇形、无 `npu_` | `complex_mul`,不要 `npu_complex_mul` / `ComplexMul` |
255-| `from __future__ import annotations` | 兼容 `X \| None` |
256- 
257-`attention_core` 这类组合封装可以放在同一模块,但必须在 [api_list.md](../api_list.md) 写明「不是独立 OPP 算子」。
258- 
259----
260- 
261-## 5. Triton-Ascend 规范
262- 
263-- 实现放 `triton-ascend/<op_family>/`,测试放 `tests/triton/<op_family>/`。
264-- **必须**提供非 NPU / 无 Triton 时的 PyTorch 回退,接口语义与 NPU 路径一致。
265-- 在 `setup.py` 登记为 `oil_gas_ops_prospect.<op_family>`,不要加 `triton` 二级包。
266-- 新 kernel 分支必须能被 `triton-ascend/tools/aot_compile_check.py` 覆盖。
267-- 约束与数学定义写入 [op_list.md](../op_list.md),签名写入 [api_list.md](../api_list.md)。
268-- 裸 checkout 由 `scripts/local_pkg.py` 挂导入名;验 whl 时置 `OIL_GAS_OPS_WHL_TEST=1`。
269- 
270-步骤见[开发指南 · Triton](operator_development_guide.md#二triton-ascend-算子)。
271- 
272----
273- 
274-## 6. 构建与打包
275- 
276-### 6.1 流程
277- 
278-```text
279-msopgen gen (DummyOp IR) → 与当前 CANN/SOC 匹配的脚手架
280- → merge_operators.py(按 --ops / manifest 合并)
281- → patch_*.py(CMakePresets / vendor / include)
282- → CustomOp/build.sh → OPP .run
283- → output/<SOC>/oil_gas_ops_prospect_opp_<os>_<arch>.run
284-```
285- 
286-顶层 `bash build.sh` 接着:安装 `.run` → 编 whl(含 pybind)→ 安装 whl → 冒烟。
287- 
288-### 6.2 添加 AscendC 算子后要做的事
289- 
290-1. `operators/<name>/op_host/` + `op_kernel/`(§2.2)。
291-2. `manifest.tsv` 追加一行。
292-3. `ascendc/pybind/<family>.py` + `csrc/<family>/bindings.cpp`。
293-4. ophost UT:改 `tests/ascendc/ophost/CMakeLists.txt`。
294-5. NPU 对比:新建 `tests/ascendc/<name>/`。
295-6. 更新 [op_list.md](../op_list.md)、[api_list.md](../api_list.md)、`ascendc/README.md`。
296-7. `bash ascendc/build.sh --pkg --ops <name>` 与 `bash build.sh -u`。
297- 
298-完整选项见 [build.md](../context/build.md)。
299- 
300-### 6.3 安装与扩展
301- 
302-```bash
303-# OPP(由顶层 build.sh 或手动安装)
304-bash output/Ascend910B1/oil_gas_ops_prospect_opp_<os>_<arch>.run --quiet
305- 
306-export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${LD_LIBRARY_PATH}
307- 
308-# 就地编 pybind(NPU 测试用,不装 whl)
309-python3 ascendc/pybind_setup.py build_ext --inplace
310-```
311- 
312-`OIL_GAS_OPS_SKIP_EXT=1` 可跳过扩展编译。CANN 根目录默认 `/usr/local/Ascend/cann-9.0.0`,用 `ASCEND_HOME_PATH` 覆盖。
313- 
314----
315- 
316-## 7. 测试与精度
317- 
318-### 7.1 分层(本仓库实际结构)
319- 
320-```text
321-无卡(CI JOB_llt / bash build.sh -u)
322- ├─ tests/ascendc/ophost/ googletest:真实 tiling_core
323- ├─ tests/triton/<family>/ pytest:CPU 回退精度
324- └─ triton-ascend/tools/aot_compile_check.py
325- 
326-有卡(PR 描述必须附结果,门禁不代跑)
327- ├─ tests/ascendc/<op>/run_test.sh
328- └─ tests/triton/<family>/verify_*.py(如有)
329-```
330- 
331-不要在 `operators/<op>/` 下再建 `testplan/`、`teste2e/`、`delivery/`。
332- 
333-### 7.2 ophost UT 必须覆盖
334- 
335-- UB 预算与 tile 上限、32B 对齐
336-- 小 tensor 收缩、空张量、`aivNum=0`、非法 shape / dtype
337-- 核间切分不变量:每个元素恰好被一个核处理一次,无空洞无重叠
338-- 融合算子:bias 广播 / 全尺寸、`Q≠S`、VEC / MIX 路径分发
339- 
340-### 7.3 NPU 对比测试
341- 
342-每个 AscendC 算子(或前反向一对)在 `tests/ascendc/<op>/`:
343- 
344-```text
345-tests/ascendc/<op>/
346-├── test_<op>_npu.py
347-├── run_test.sh # 安装 OPP → 就地编扩展 → 跑脚本
348-└── README.md
349-```
350- 
351-基准必须是等价的 PyTorch / torch_npu 表达式。记录相对误差或契约式 `atol/rtol`、时延与加速比、是否原地写回。
352- 
353-### 7.4 精度标准
354- 
355-以[生态算子开源精度标准](https://gitcode.com/cann/opbase/blob/master/docs/zh/ops_precision_standard/experimental_standard.md)为准。本仓已交付算子的验收阈值:
356- 
357-| 算子 | dtype | 阈值 | 出处 |
358-|------|-------|------|------|
359-| `ComplexMul` / `Grad` | fp32 | 相对误差默认 `1e-5` | `tests/ascendc/complex_mul/` |
360-| `FusedBiasSoftmax` / `FusedSoftmaxGrad` | fp32 | `max_diff ≤ 1e-5` | `tests/ascendc/fused_bias_softmax/` |
361-| 同上 | bf16 | `atol=2e-2, rtol=4e-3` | 同上(量化半 ulp 随幅值增长,不用纯 abs) |
362- 
363-新算子在 PR 中按 [精度验收报告模板](precision_acceptance_template.md) 填写,通过率 100%。严重 / 主要遗留问题未闭环不得合入。
364- 
365-### 7.5 真机性能与内存(推荐)
366- 
367-门禁不跑 NPU。合入前应在 910B 上:
368- 
369-- 精度脚本:预热后多次迭代,`torch.npu.synchronize()`,记 `ms/call`。
370-- 需要定位瓶颈时用 `msprof`(见 [op_debug_prof.md](../debug/op_debug_prof.md))。
371-- 怀疑越界 / 踩踏时用 `mssanitizer --tool=memcheck --leak-check=no --check-device-heap=no`(PyTorch 缓存分配器会误报 leak / device heap)。
372- 
373-### 7.6 文档同步
374- 
375-算子语义或约束变了,必须同步 [op_list.md](../op_list.md)、[api_list.md](../api_list.md)、对应 `README.md`。
376- 
377----
378- 
379-## 8. PR 产物清单
380- 
381-### 8.1 必须提交
382- 
383-| 类别 | 路径 |
384-|------|------|
385-| Host / Kernel | `ascendc/operators/<op>/op_host/`、`op_kernel/` |
386-| 注册表 | `ascendc/operators/manifest.tsv` 新行 |
387-| Python API | `ascendc/pybind/<family>.py` |
388-| aclnn 绑定 | `ascendc/pybind/csrc/<family>/bindings.cpp` |
389-| ophost UT | `tests/ascendc/ophost/test_<op>_tiling.cpp` + `CMakeLists.txt` |
390-| NPU 对比 | `tests/ascendc/<op>/`(Triton 则为 `tests/triton/<family>/`) |
391-| 文档 | `docs/zh/op_list.md`、`api_list.md`;AscendC 还要 `ascendc/README.md` |
392-| 精度报告 | PR 描述附 [模板](precision_acceptance_template.md) |
393-| Triton 登记 | `setup.py` 的 `packages` / `package_dir`;AOT 检查覆盖 |
394- 
395-不要求、也不要新增:算子目录内的 `delivery/` SHA256 副本、`testplan/`、`teste2e/`。
396- 
397-### 8.2 不提交
398- 
399-`output/`、`ascendc/_work/`、`build/`、`*.egg-info/`、`*.so`、`__pycache__/`、`.coverage`、`mindstudio_sanitizer_log/`。已在 `.gitignore`。
400- 
401----
402- 
403-## 9. 审核与迭代
404- 
405-| 优先级 | 含义 | 处理 |
406-|--------|------|------|
407-| **P0** | 功能错误、精度超限、crash、梯度断裂 | 阻塞合入 |
408-| **P1** | 切分不可测、命名违规、静默回退掩盖失败 | 必须修 |
409-| **P2** | 文档不同步、缺边界用例 | 应修 |
410-| **P3** | 性能建议 | 可选 |
411- 
412-流程按 [CONTRIBUTING.md](../../../CONTRIBUTING.md):先 Issue,再 PR;门禁 `compile`;Committer `/lgtm`,Maintainer `/approve`。
413- 
414----
415- 
416-## 10. 分支与提交
417- 
418-| 类型 | 格式 | 示例 |
419-|------|------|------|
420-| 功能 | `feat/<op-name>` | `feat/complex-mul` |
421-| 修复 | `fix/<op-name>` | `fix/fused-softmax-grad-tiling` |
422-| 文档 | `docs/<topic>` | `docs/operator-dev-spec` |
423-| 布局 / 重构 | `refactor/<topic>` | `refactor/opensource-layout` |
424- 
425-Conventional Commits:`feat` / `fix` / `test` / `docs` / `refactor` / `chore`。描述可中文。
426- 
427-PR 标题写清算子或主题;描述含关联 Issue、测试统计、精度报告。CLA 邮箱与 `git config user.email` 一致。
428- 
429----
430- 
431-## 11. 附录:检查清单
432- 
433-```
434-□ operators/<op>/op_host + op_kernel(含 tiling_core.h)
435-□ manifest.tsv 新行;list-ops 能看到
436-□ pybind/<family>.py:蛇形公开名、autograd、torch 参考、OIL_GAS_OPS_* 开关
437-□ pybind/csrc/<family>/bindings.cpp(直调 aclnn)
438-□ tests/ascendc/ophost 已接入;bash build.sh -u 通过
439-□ tests/ascendc/<op>/ 或 tests/triton/<family>/
440-□ op_list.md / api_list.md / 相关 README 已更新
441-□ 许可证头(Apache License 2.0)
442-□ bash scripts/check_license_header.sh
443-□ 910B 真机精度 + 性能写入 PR 模板;通过率 100%
444-□ 无卡门禁(compile / 静态检查 / UT)通过
445-```
446- 
447-### 环境变量速查
448- 
449-| 变量 | 默认 | 作用 |
450-|------|------|------|
451-| `ASCEND_HOME_PATH` | `/usr/local/Ascend/cann-9.0.0` | CANN 根目录 |
452-| `LD_LIBRARY_PATH` | — | 需含 `opp/vendors/customize/op_api/lib` |
453-| `OIL_GAS_OPS_<OP>_IMPL` | `torch` | `torch` / `ascendc` / `auto` |
454-| `OIL_GAS_OPS_REQUIRE_ASCENDC_<OP>` | `0` | `1` 禁止回退 |
455-| `OIL_GAS_OPS_OUTPUT_DIR` | `<repo>/output` | 产物目录 |
456-| `OIL_GAS_OPS_SKIP_EXT` | `0` | `1` 跳过 pybind 扩展 |
457-| `OIL_GAS_OPS_WHL_TEST` | `0` | `1` 时禁止导入仓内源码 |
458-| `OIL_GAS_OPS_RUN_PKG` | 自动探测 | NPU 测试指定 `.run` 包 |
459- 
460-完整列表见 [api_list.md · 环境变量](../api_list.md#环境变量)。
@@ -1,6 +1,6 @@
1# 算子开发指南1# 算子开发指南
2 2 
3-本文讲**怎么做**(目录、注册、编译、测试)。必须遵守的命名、切分可测、Python 封装、PR 交付件与精度验收见[算子开发规格](operator-dev-spec.md)。3+本文说明如何新增 AscendC / Triton-Ascend 算子,以及必须遵守的目录、命名、封装与验收约定。路径以 [`dir_structure.md`](../context/dir_structure.md) 为准。Issue / PR 流程见[贡献指南](../../../CONTRIBUTING.md)。
4 4 
5本仓有两类算子后端,交付件与开发流程不同。请先判断新算子属于哪一类:5本仓有两类算子后端,交付件与开发流程不同。请先判断新算子属于哪一类:
6 6 
@@ -28,11 +28,13 @@ ascendc/operators/${op_name}/
28 └── ${op_name}_tiling.h # kernel 侧 TilingData 视图28 └── ${op_name}_tiling.h # kernel 侧 TilingData 视图
29```29```
30 30 
31-多算子共用的 host / kernel 头放在 `ascendc/operators/_common/`,合并进 msopgen 工程时会平铺到 `op_host/` 与 `op_kernel/`。31+`operators/` 只放实现。C++ 测试与样例放在仓库根 `tests/`、`examples/`。多算子共用的 host / kernel 头放在 `ascendc/operators/_common/`,合并进 msopgen 工程时会平铺到 `op_host/` 与 `op_kernel/`。
32+ 
33+一对前反向可以共用一个 Python 模块,但 **manifest 与 `operators/` 各占一行 / 一个目录**。共享逻辑放 `_common/`。
32 34 
33### 2. 注册算子35### 2. 注册算子
34 36 
35-在 `ascendc/operators/manifest.tsv` 追加一行(**tab 分隔**)。可用如下命令确认注册成功:37+在 `ascendc/operators/manifest.tsv` 追加一行(**tab 分隔**:目录名、OpType、kernel 入口文件名)。可用如下命令确认注册成功:
36 38 
37```bash39```bash
38cd ascendc && bash build.sh list-ops40cd ascendc && bash build.sh list-ops
@@ -40,48 +42,87 @@ cd ascendc && bash build.sh list-ops
40 42 
41### 3. 切分算法必须可测43### 3. 切分算法必须可测
42 44 
43-`<Op>TilingFunc` / `InferShape` 的入口签名依赖 `gert::TilingContext`,`Compute<Op>Tiling` 依赖 `platform_ascendc::PlatformAscendC`;CANN 二进制包未提供 gert faker(该基建在 metadef 源码仓),两者都无法直接 UT。45+`<Op>TilingFunc` / `InferShape` 的入口签名依赖 `gert::TilingContext`,`Compute<Op>Tiling` 依赖 `platform_ascendc::PlatformAscendC`;CANN 二进制包未提供 gert faker,两者都无法直接 UT。
44 46 
45因此约定:**把与平台无关的切分算法抽到 `op_host/<op>_tiling_core.h`**,UT 直接驱动核心函数,覆盖的是真实算法。47因此约定:**把与平台无关的切分算法抽到 `op_host/<op>_tiling_core.h`**,UT 直接驱动核心函数,覆盖的是真实算法。
46 48 
49+Host 侧用 `BEGIN_TILING_DATA_DEF`(opbuild 需要);kernel 侧用 **plain struct** 镜像同一布局(`ccec` 不识别该宏)。`TilingFunc` 负责:形状 / dtype 校验 → 调 `*_tiling_core.h` → `SetBlockDim` / `SetTilingKey` → `SaveToBuffer`。`GetCoreNumAiv()` 为 0 必须失败或给出明确兜底,并在 ophost UT 覆盖该分支。
50+ 
47现有算子的切分策略可作参考:按元素总数在 AIV 核间均分(32B 对齐),核内再按 UB 预算分拍。51现有算子的切分策略可作参考:按元素总数在 AIV 核间均分(32B 对齐),核内再按 UB 预算分拍。
48 52 
49### 4. 命名约定53### 4. 命名约定
50 54 
51-- **C++ / CANN**:OpType 用 PascalCase(如 `ComplexMul`),生成的 aclnn 接口为 `aclnn` + OpType(如 `aclnnComplexMul`)。55+| 侧 | 规则 | 示例 |
52-- **Python**:公开函数用蛇形命名(如 `complex_mul`),不要用驼峰,也不要加 `npu_` 前缀。56+|----|------|------|
53-- kernel 入口函数名通过算子 def 的 `opFile.value` / `opInterface.value` 固定为蛇形(如 `complex_mul`),OpType 改名不影响 kernel 源码;aclnn 接口名随 OpType 变化。57+| C++ OpType | PascalCase | `ComplexMul`、`FusedBiasSoftmax` |
58+| aclnn | `aclnn` + OpType | `aclnnComplexMul` |
59+| kernel 入口 / `opFile.value` | 蛇形,与源文件名一致 | `complex_mul` |
60+| Python 公开函数 | 蛇形,**不加** `npu_` 前缀 | `complex_mul`、`fused_bias_softmax` |
61+| 反向 OPP 算子 | `Grad` 后缀 | `ComplexMulGrad`、`FusedSoftmaxGrad` |
62+| 环境变量 | `OIL_GAS_OPS_` 前缀 | `OIL_GAS_OPS_COMPLEX_MUL_IMPL` |
54 63 
55-### 5. 补充单元测试64+OpType 改名不影响 kernel 源码(入口由 `opFile.value` / `opInterface.value` 固定),但 aclnn 符号随 OpType 变化。
56 65 
57-测试用例统一收在仓库根目录的 `tests/` 下,算子目录只放实现。把新算子的 `op_host` 目录加进 `tests/ascendc/ophost/CMakeLists.txt` 的 `OPHOST_UT_INCLUDE_DIRS`,并把测试源加进 `ophost_ut` 目标:66+### 5. OpDef 与 kernel 入口
58 67 
59-```cmake68+使用 `ops::OpDef` + `OP_ADD`,不要用旧的 `REG_OP`。`opFile.value` / `opInterface.value` 必须等于 `op_kernel/` 下入口文件名(无扩展名)。需要 aclnn 时设 `aclnnSupport.value = support_aclnn`(或以该算子实测为准)。逐元素 float32/ND 算子优先复用 `_common/op_host/og_tiling_context.h` 的 `DeclareRequiredFloatNdIo` / `AddAscend910bElementwiseConfig`。
60-set(OPHOST_UT_INCLUDE_DIRS
61- "${OPERATORS_DIR}/_common/op_host"
62- "${OPERATORS_DIR}/complex_mul/op_host"
63- "${OPERATORS_DIR}/${op_name}/op_host" # 新增
64-)
65 69 
66-add_executable(ophost_ut70+入口符号必须与 `opFile.value` 一致。Ascend 910B 向量指令以 **64 个元素**为粒度:`count` 不是 64 倍数时,尾部可能被跳过或写成 0。计算用 aligned 长度,`DataCopyPad` 写出用原始 `count`,禁止越界写。UB 按 32B 规划。
67- test_complex_mul_tiling.cpp71+ 
68- test_${op_name}_tiling.cpp # 新增72+### 6. Python 封装
69-)73+ 
74+| 文件 | 作用 |
75+|------|------|
76+| `ascendc/pybind/<op_family>.py` | 公开 API,打包为 `oil_gas_ops_prospect.<op_family>` |
77+| `ascendc/pybind/csrc/<op_family>/bindings.cpp` | 直调 aclnn 两段式入口,编成 `<op_family>_npu` |
78+ 
79+新建 `csrc/<op_family>/*.cpp` 后,`pybind_extensions.py` 会自动带上扩展,不必再手写 `setup.py` 的 `CppExtension`。写法以现有 `complex_mul.py` 为准:
80+ 
81+| 规则 | 说明 |
82+|------|------|
83+| 函数式 API 走 `autograd.Function.apply` | 直接调 `ext.forward()` 会断梯度图 |
84+| 必须有独立 PyTorch 参考 | 默认 `IMPL=torch`;对拍依赖它 |
85+| 实现开关 | `OIL_GAS_OPS_<OP>_IMPL=torch\|ascendc\|auto`,默认 `torch` |
86+| 强制自定义算子 | `OIL_GAS_OPS_REQUIRE_ASCENDC_<OP>=1` 时禁止静默回退 |
87+| `@lru_cache` 缓存扩展 | 避免重复 import |
88+| 公开名蛇形、无 `npu_` | `complex_mul`,不要 `npu_complex_mul` / `ComplexMul` |
89+ 
90+`attention_core` 这类组合封装可以放在同一模块,但必须在 [api_list.md](../api_list.md) 写明「不是独立 OPP 算子」。
91+ 
92+### 7. 补充单元测试
93+ 
94+AscendC 的 C++ 测试放在仓库根,算子目录只保留实现。本仓没有 ATK / `OP_API_UT`。
95+ 
96+```text
97+tests/ut/op_host/${op_name}/test_${op_name}_tiling.cpp # tiling UT(googletest)
98+tests/st/aclnn<OpType>/ # 多 case 两段式 aclnn + 真值(需要 NPU)
99+examples/aclnn/${op_name}/test_aclnn_${op_name}.cpp # 内联小 tensor,两段式 aclnn(需要 NPU)
70```100```
71 101 
72-建议覆盖:UB 预算约束、对齐与上限、小 tensor 收缩、退化输入(负元素数、`aivNum=0`、空张量),以及核间切分的覆盖性不变量(每个元素恰好被一个 AIV 核处理一次、区间无空洞无重叠)。102+`tests/ascendc/ophost/CMakeLists.txt` 会自动收集 `tests/ut/op_host/*/test_*.cpp`
103+并加上各 `op_host` include。Python 精度对比放在 `tests/ascendc/${op_name}/`。
104+登记 ST / 精度 / 样例到 `tests/suites.json`。
73 105 
74-### 6. 编译与验证106+ophost UT 建议覆盖:UB 预算约束、对齐与上限、小 tensor 收缩、退化输入(负元素数、`aivNum=0`、空张量),以及核间切分的覆盖性不变量(每个元素恰好被一个 AIV 核处理一次、区间无空洞无重叠)。融合算子另含 bias 广播 / 全尺寸、`Q≠S`、VEC / MIX 路径分发。
107+ 
108+### 8. 编译与验证
75 109 
76```bash110```bash
77-# 打包111+# 打包(AscendC 侧只出 .run)
78cd ascendc && bash build.sh --pkg --ops ${op_name}112cd ascendc && bash build.sh --pkg --ops ${op_name}
79 113 
80-# 无卡 UT114+# 仓库根目录
81-bash build.sh -u --ophost # 仓库根目录;或 cd ascendc && bash build.sh -u --ophost115+bash build.sh --pkg # 只出 OPP .run
116+bash build.sh # 编 .run 与 whl,不安装
117+bash build.sh --install # 编译并安装
118+ 
119+# 910B 上板:UT + ST + 精度 + 性能(写出 test_result/test_report.md)
120+bash build.sh -u --ut --st --precision --perf --ops=${op_name}
82```121```
83 122 
84-### 7. NPU 精度/性能测试123+完整选项见 [build.md](../context/build.md)。
124+ 
125+### 9. NPU 精度/性能测试
85 126 
86在 `tests/ascendc/${op_name}/` 下按现有结构组织,统一以等价的 PyTorch/torch_npu 表达式为基准:127在 `tests/ascendc/${op_name}/` 下按现有结构组织,统一以等价的 PyTorch/torch_npu 表达式为基准:
87 128 
@@ -99,7 +140,17 @@ tests/ascendc/${op_name}/
99`ascendc/pybind/csrc/<op_family>/` 下的 `*.cpp` 会自动编译成扩展模块 `<op_family>_npu`,140`ascendc/pybind/csrc/<op_family>/` 下的 `*.cpp` 会自动编译成扩展模块 `<op_family>_npu`,
100无需改构建脚本,详见 [`ascendc/pybind/README.md`](../../../ascendc/pybind/README.md)。141无需改构建脚本,详见 [`ascendc/pybind/README.md`](../../../ascendc/pybind/README.md)。
101 142 
102-### 8. 移植 CANN 框架算子的注意事项143+精度标准以[生态算子开源精度标准](https://gitcode.com/cann/opbase/blob/master/docs/zh/ops_precision_standard/experimental_standard.md)为准。本仓已交付算子的验收阈值:
144+ 
145+| 算子 | dtype | 阈值 | 出处 |
146+|------|-------|------|------|
147+| `ComplexMul` / `Grad` | fp32 | 相对误差默认 `1e-5` | `tests/ascendc/complex_mul/` |
148+| `FusedBiasSoftmax` / `FusedSoftmaxGrad` | fp32 | `max_diff ≤ 1e-5` | `tests/ascendc/fused_bias_softmax/` |
149+| 同上 | bf16 | `atol=2e-2, rtol=4e-3` | 同上 |
150+ 
151+新算子在 PR 中按 [精度验收报告模板](precision_acceptance_template.md) 填写,通过率 100%。严重 / 主要遗留问题未闭环不得合入。算子在 NPU 上的精度、性能、显存必须在真实 910B 上验证通过。
152+ 
153+### 10. 移植 CANN 框架算子的注意事项
103 154 
104部分从 CANN 原仓移植的算子的 host tiling/infershape 依赖 CANN 内部框架头(`op_host/tiling_base.h`、`op_host/tiling_templates_registry.h`、`err/ops_err.h`、`log/log.h` 等),这些头在 msopgen 自定义算子工程的 opbuild 编译环境下缺失或会拉入 slog 依赖。处理方式(保持源码改动最小):155部分从 CANN 原仓移植的算子的 host tiling/infershape 依赖 CANN 内部框架头(`op_host/tiling_base.h`、`op_host/tiling_templates_registry.h`、`err/ops_err.h`、`log/log.h` 等),这些头在 msopgen 自定义算子工程的 opbuild 编译环境下缺失或会拉入 slog 依赖。处理方式(保持源码改动最小):
105 156 
@@ -109,6 +160,8 @@ tests/ascendc/${op_name}/
109 160 
110更多细节见 [`ascendc/README.md`](../../../ascendc/README.md)。161更多细节见 [`ascendc/README.md`](../../../ascendc/README.md)。
111 162 
163+当前交付与真机验证面向 **Ascend 910B**(默认 `Ascend910B1`)。新增 SOC 必须同时写入 `SUPPORTED_SOC_VERSIONS`、在 OpDef 的 `AICore().AddConfig(...)` 中增加对应配置,并在该 SOC 上完成精度验收。
164+ 
112---165---
113 166 
114## 二、Triton-Ascend 算子167## 二、Triton-Ascend 算子
@@ -119,20 +172,22 @@ tests/ascendc/${op_name}/
119triton-ascend/${op_family}/172triton-ascend/${op_family}/
120├── __init__.py # 对外导出173├── __init__.py # 对外导出
121├── triton_${op_family}.py # Triton kernel 与 autograd 封装174├── triton_${op_family}.py # Triton kernel 与 autograd 封装
122-├── test_triton_${op_family}.py # CPU 回退精度测试(无卡可跑)
123└── README.md # 算子说明175└── README.md # 算子说明
176+ 
177+tests/triton/${op_family}/
178+└── test_triton_${op_family}.py # CPU 回退路径精度测试
124```179```
125 180 
126### 2. 必须提供 CPU 回退181### 2. 必须提供 CPU 回退
127 182 
128-kernel 需提供**非 NPU 环境下的 PyTorch 回退实现**,保证无卡环境可做精度自检。现有 `real_rfft` 的写法可作参考:183+kernel 需提供 **PyTorch 回退实现**,接口语义与 NPU 路径一致。现有 `real_rfft` 的写法可作参考:
129 184 
130```python185```python
131def real_rfft(x, dim, norm="ortho", ...):186def real_rfft(x, dim, norm="ortho", ...):
132 if norm != "ortho":187 if norm != "ortho":
133 raise ValueError(f"only norm='ortho' is supported, got {norm!r}")188 raise ValueError(f"only norm='ortho' is supported, got {norm!r}")
134 if not (_is_npu_tensor(x) and _has_triton()):189 if not (_is_npu_tensor(x) and _has_triton()):
135- return _torch_rfft_real(x, dim=dim, norm=norm) # CPU / 无 Triton 回退190+ return _torch_rfft_real(x, dim=dim, norm=norm)
136 return _RealRFFTFunction.apply(x, dim, norm, ...)191 return _RealRFFTFunction.apply(x, dim, norm, ...)
137```192```
138 193 
@@ -151,9 +206,13 @@ package_dir={
151},206},
152```207```
153 208 
209+不要加 `triton` / `ascendc` 二级包。裸 checkout 由 `scripts/local_pkg.py` 挂导入名;验 whl 时置 `OIL_GAS_OPS_WHL_TEST=1`。
210+ 
154### 4. 纳入 AOT 编译检查211### 4. 纳入 AOT 编译检查
155 212 
156-新增 kernel 分支需能被 `triton-ascend/tools/aot_compile_check.py` 覆盖,保证无卡环境也能验证全部分支可编译(TTIR → Linalg → bishengir → npubin)。213+新增 kernel 分支需能被 `triton-ascend/tools/aot_compile_check.py` 覆盖,保证全部分支可编译(TTIR → Linalg → bishengir → npubin)。
214+ 
215+约束与数学定义写入 [op_list.md](../op_list.md),签名写入 [api_list.md](../api_list.md)。
157 216 
158### 5. 本地源码 / whl 切换217### 5. 本地源码 / whl 切换
159 218 
@@ -175,19 +234,26 @@ bash tests/run_whl_tests.sh
175## 三、提交前自检234## 三、提交前自检
176 235 
177```bash236```bash
178-# 全量无卡单元测试237+bash build.sh --install
179bash build.sh -u238bash build.sh -u
180- 
181-# 许可证头检查
182bash scripts/check_license_header.sh239bash scripts/check_license_header.sh
183- 
184-# 提交前钩子(可选)
185-pip install pre-commit && pre-commit install
186-pre-commit run --all-files
187```240```
188 241 
189-新增源文件需带 Apache License 2.0 许可证头,格式参考仓内已有文件。242+新增源文件需带 Apache License 2.0 许可证头,格式参考仓内已有文件。算子语义或约束变了,必须同步 [op_list.md](../op_list.md)、[api_list.md](../api_list.md)、对应 `README.md`。
190 243 
191-> **编译通过不等于功能正确。** 算子在 NPU 上的精度、性能、显存必须在真实 910B 上验证,并在 PR 描述中附上[精度验收报告](precision_acceptance_template.md)。244+### 交付件清单
192 245 
193-完整的 Issue / PR 流程见[贡献指南](../../../CONTRIBUTING.md);提交前按[规格附录检查清单](operator-dev-spec.md#11-附录检查清单)自检。246+| 类别 | 路径 |
247+|------|------|
248+| Host / Kernel | `ascendc/operators/<op>/op_host/`、`op_kernel/`(含 `*_tiling_core.h`) |
249+| 注册表 | `ascendc/operators/manifest.tsv` 新行 |
250+| Python API | `ascendc/pybind/<family>.py`(蛇形公开名、autograd、torch 参考、`OIL_GAS_OPS_*` 开关) |
251+| aclnn 绑定 | `ascendc/pybind/csrc/<family>/bindings.cpp` |
252+| ophost UT | `tests/ut/op_host/<op>/test_<op>_tiling.cpp` |
253+| aclnn example / ST | `examples/aclnn/<op>/`、`tests/st/aclnn<OpType>/`;登记 `tests/suites.json` |
254+| NPU 对比 | `tests/ascendc/<op>/`(Triton 则为 `tests/triton/<family>/`) |
255+| 文档 | `docs/zh/op_list.md`、`api_list.md`;AscendC 还要 `ascendc/README.md` |
256+| 精度报告 | PR 描述附 [模板](precision_acceptance_template.md) |
257+| Triton 登记 | `setup.py` 的 `packages` / `package_dir`;AOT 检查覆盖 |
258+ 
259+不要在 `operators/<op>/` 下再建 `tests/`、`examples/`、`testplan/`、`teste2e/`、`delivery/`。`output/`、`ascendc/_work/`、`build/`、`*.so` 等已在 `.gitignore`,不要提交。
@@ -77,10 +77,11 @@
77 77 
78| 测试项 | 命令 | 用例数 | 通过数 | 结论 |78| 测试项 | 命令 | 用例数 | 通过数 | 结论 |
79|--------|------|--------|--------|------|79|--------|------|--------|--------|------|
80-| ophost tiling UT | `bash ascendc/build.sh -u --ophost` | | | |80+| ophost tiling UT | `bash build.sh -u --ophost` | | | |
81+| ST | `bash build.sh -u --st --ops=<op>` | | | |
82+| 精度 / 性能 | `bash build.sh -u --precision --perf --ops=<op>` | | | |
81| Triton CPU 回退精度 | `python3 -m pytest tests/triton/real_fft -q` | | | |83| Triton CPU 回退精度 | `python3 -m pytest tests/triton/real_fft -q` | | | |
82| Triton AOT 编译检查 | `python3 triton-ascend/tools/aot_compile_check.py` | | | |84| Triton AOT 编译检查 | `python3 triton-ascend/tools/aot_compile_check.py` | | | |
83-| NPU 精度/性能对比 | `bash tests/ascendc/<op>/run_test.sh` | | | |
84 85 
85要求:全量 UT 通过且日志无 ERROR。86要求:全量 UT 通过且日志无 ERROR。
86 87 
@@ -9,7 +9,7 @@ source /usr/local/Ascend/cann-9.0.0/set_env.sh
9export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:$LD_LIBRARY_PATH9export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:$LD_LIBRARY_PATH
10```10```
11 11 
12-并已通过 `bash build.sh` 完成 OPP 包与 whl 的安装。12+并已通过 `bash build.sh --install` 完成 OPP 包与 whl 的安装。
13 13 
14## AscendC:复数逐点乘14## AscendC:复数逐点乘
15 15 
@@ -1,250 +0,0 @@
1-# oil-gas-ops-prospect仓开源测试报告
2- 
3-## 1. 概述
4- 
5-本报告为 oil-gas-ops-prospect 项目开源测试报告。测试基于 CANN 9.0.0 开源包,覆盖仓内已交付算子的文档、编译出包、安装卸载、功能、精度与性能,以及配套环境兼容性。测试过程中对已纳入范围的特性执行覆盖率为 100%。
6- 
7-本仓面向油气勘探计算负载,按算子名对外提供 Python / aclnn 接口,不绑定单一上层模型。
8- 
9-### 新增特性清单
10- 
11-| 序号 | 特性名称 | 特性描述 |
12-| --- | --- | --- |
13-| 1 | ComplexMul | 复数逐点乘:`(ur, ui), (vr, vi) → (wr, wi)`,`float32` / `ND`,全部张量同 shape |
14-| 2 | ComplexMulGrad | `ComplexMul` 的反向:由 `gwr, gwi` 与前向输入得到 `gur, gui, gvr, gvi`,`float32` / `ND` |
15-| 3 | FusedBiasSoftmax | `y = softmax(x + bias1 + bias2, -1)`,`bias1`/`bias2` 可选,NPU 路径原地写回 `x`,支持 `fp32` / `fp16` / `bf16` |
16-| 4 | FusedSoftmaxGrad | 注意力 softmax JVP:`g = go @ V^T` 后原地覆写 `probs`,支持 `Q≠S`,`fp32` / `fp16` / `bf16` |
17-| 5 | real_rfft | 实数正向 FFT:`x → (xr, xi)`,`norm="ortho"`,频率轴长度为 `n//2+1` |
18-| 6 | real_irfft | 由半谱实部、虚部重建实数:`(xr, xi), n → y`,`norm="ortho"` |
19- 
20-`attention_core(q, k, v, bias1, bias2)` 为 FusedBiasSoftmax 与 FusedSoftmaxGrad 按 openfold 注意力语义的组合封装,随上述融合算子一并验证。
21- 
22-### 解决问题列表
23- 
24-不涉及(首个开源测试版本,无历史问题清单)。
25- 
26-### 测试活动
27- 
28-本次测试活动包括:Readme 与算子文档评估、一键编译出包、算子包与 Python 包安装卸载、按算子的功能/精度/性能验证、可靠性与兼容性验证。
29- 
30-## 2. 版本测试信息
31- 
32-本节描述被测对象的版本信息和测试时间,包括依赖环境与测试源码仓。
33- 
34-**硬件和版本要求**
35- 
36-产品型号:Atlas A2 / Ascend 910B(文档默认 `SOC_VERSION=Ascend910B1`)
37- 
38-操作系统:覆盖以下测试因子组合
39- 
40-| | 因子名称 | 因子取值 |
41-| --- | --- | --- |
42-| 1 | OS | Ubuntu / openEuler |
43-| 2 | CPU 架构 | ARM64 |
44-| 3 | 形态 | docker / host+device |
45- 
46-CANN 版本:CANN 9.0.0
47- 
48-驱动版本:Ascend HDK 26.0.RC1
49- 
50-Python 版本:Python >= 3.10(本次实测 3.11)
51- 
52-PyTorch 版本:2.7.1(配套 torch_npu 2.7.1)
53- 
54-依赖三方库版本:与 CANN 9.0.0 配套的 torch_npu、Triton-Ascend 3.2.1;构建依赖见仓库 `requirements.txt`
55- 
56-测试源码仓:https://gitcode.com/cann/oil-gas-ops-prospect
57- 
58-测试日期:2026-09-12
59- 
60-## 3. 测试结论
61- 
62-给出明确的版本发布结论。
63- 
64-oil-gas-ops-prospect 基于 CANN 9.0.0 完成开源测试,共计执行 18 组特性级测试用例,发现 0 个问题,其中有效问题 0 个。整体质量良好,满足出口质量标准。**建议发布。**
65- 
66-| 评估维度 | 评估结果 |
67-| --- | --- |
68-| 测试用例总数 | 18 组(按已交付算子与关键测试活动归并,不含参数化子项展开) |
69-| 测试执行通过率 | 100%(18/18) |
70-| 测试结论 | 通过 |
71- 
72-## 4. 特性质量评估
73- 
74-新特性的验证结果清单如下。功能、精度在 CPU 参考路径与 NPU 路径上一并验证(非 NPU 环境走 PyTorch 回退,NPU 环境与同设备标杆对拍)。
75- 
76-| 序号 | 特性 | 测试结论 | 功能 | 精度 | 性能 | 可靠性 | 兼容性 |
77-| --- | --- | --- | --- | --- | --- | --- | --- |
78-| 1 | ComplexMul | 通过 | Pass | Pass | Pass | Pass | Pass |
79-| 2 | ComplexMulGrad | 通过 | Pass | Pass | Pass | Pass | Pass |
80-| 3 | FusedBiasSoftmax | 通过 | Pass | Pass | Pass | Pass | Pass |
81-| 4 | FusedSoftmaxGrad | 通过 | Pass | Pass | Pass | Pass | Pass |
82-| 5 | real_rfft | 通过 | Pass | Pass | Pass | Pass | Pass |
83-| 6 | real_irfft | 通过 | Pass | Pass | Pass | Pass | Pass |
84- 
85-特性级用例(按特性介绍组织,同一用例同时覆盖非 NPU 回退与 NPU 对拍,不按实现后端拆分):
86- 
87-| 序号 | 测试用例 | 测试内容 | 结果 |
88-| --- | --- | --- | --- |
89-| 1 | ComplexMul 切分与约束 | 代表 shape / 小 tensor / 退化输入的切分计划,核间覆盖与对齐,UB 预算 | PASS |
90-| 2 | ComplexMul 前反向精度 | 闭式解对拍(含 `whl` 接口与 NPU 执行):`(4,1024)`、`(1,1,192,320,320)` 等,相对误差阈值 `1e-5` | PASS |
91-| 3 | ComplexMulGrad 切分与梯度 | 反向 tile 小于前向、代表 shape 填满、退化输入;梯度与 autograd / 闭式解一致 | PASS |
92-| 4 | FusedBiasSoftmax 切分与约束 | 双 bias 广播、无 bias、非对齐宽、满尺寸 bias1、非法输入拒绝、dtype 对应 tiling key | PASS |
93-| 5 | FusedBiasSoftmax 前向精度 | `softmax(x+bias1+bias2,-1)`:bf16 广播/全尺寸/尖峰分布、fp32 无 bias;原地语义;同 shape 不同缓冲互不污染;与 fp32 标杆对拍 | PASS |
94-| 6 | FusedSoftmaxGrad 切分与路径 | 对齐 MIX、非对齐 VEC、fp32/fp16 走向量、超 UB 拒绝、`Q≠S` 合法、非法输入拒绝 | PASS |
95-| 7 | FusedSoftmaxGrad 反向精度 | `g=go@V^T` 后 JVP 原地覆写:典型 / 超核 / 非对齐 D / fp32 / `Q≠S`;与 fp32 标杆对拍 | PASS |
96-| 8 | 融合注意力链路 | `attention_core` 前向输出与 `dL/dlogits`、`dL/dq`、`dL/dv` 对拍,容差 `max_diff<0.6` | PASS |
97-| 9 | real_rfft 基函数与前向 | Fourier 基与 `torch.fft` 一致;多 shape/dim 前向半谱对拍(含回退路径与 NPU) | PASS |
98-| 10 | real_irfft 重建与往返 | 半谱重建实信号,往返误差与 `torch.fft.irfft(..., norm="ortho")` 对齐 | PASS |
99-| 11 | real_rfft / real_irfft 反向 | 多 shape/dim 反向与 CPU `torch.fft` 对拍 | PASS |
100-| 12 | real_rfft / real_irfft 大形态 | 三维潜变量形态与 `128^3` 往返/反向;默认示例可运行 | PASS |
101-| 13 | 算子文档 | 算子清单、数学定义、约束、接口列表与实现一致;快速入门与开发指南可跟随 | PASS |
102-| 14 | 编译出包 | `bash build.sh --help` 合法选项可打印;`--soc` / `--ops` / `-j` 可组合;产出 OPP `.run` 与 Python whl | PASS |
103-| 15 | 安装与卸载 | OPP 写入 customize vendor;`pip install` / `pip uninstall oil-gas-ops-prospect` 可执行 | PASS |
104-| 16 | 公开接口冒烟 | `complex_mul` / `fused_bias_softmax` / `fused_softmax_grad` / `real_rfft` / `real_irfft` 可导入;CPU 回退前反向可跑通 | PASS |
105-| 17 | 可靠性 | 非法输入被切分侧拒绝;融合算子保持原地写回;非 NPU 自动回退且接口语义不变 | PASS |
106-| 18 | 配套兼容 | CANN 9.0.0 + 配套 PyTorch / torch_npu 可完成编译、安装与算子验证 | PASS |
107- 
108-## 5. DFX专项质量评估
109- 
110-### 5.1 安全测试
111- 
112-版本整体安全测试结论:满足开源发布对安全声明与许可证的要求。
113- 
114-仓内提供 `SECURITY.md`、Apache 2.0 `LICENSE`;构建入口对非法参数拦截并提示。本次未开展独立漏洞扫描、fuzz 与交付件病毒扫描。
115- 
116-| 序号 | 测试项 | 测试结论 | 遗留风险 |
117-| --- | --- | --- | --- |
118-| 1 | 安全声明与许可证 | 通过 | 无 |
119-| 2 | 构建参数校验 | 通过 | 无 |
120- 
121-### 5.2 可靠性测试
122- 
123-可靠性指标的验证达成情况如下。
124- 
125-| 序号 | 可靠性特性 | 测试结论 | 遗留风险 |
126-| --- | --- | --- | --- |
127-| 1 | 非法或退化输入被拒绝或收缩为合法切分 | 通过 | 无 |
128-| 2 | 融合算子原地写回(输出与输入共享 device 指针) | 通过 | 无 |
129-| 3 | 非 NPU 或扩展不可用时回退 PyTorch,接口不变 | 通过 | 无 |
130-| 4 | 同 shape 重复调用精度稳定(含不同缓冲、融合注意力链路) | 通过 | 无 |
131- 
132-### 5.3 性能测试
133- 
134-该版本的性能按算子微基准测量:预热 5 次、计时 20 次取均值,加速比 = 同设备 Torch 小算子时延 / 本算子时延。测试环境为 Atlas A2 / CANN 9.0.0,不绑定具体主机。四类算子均于 2026-09-13 对同设备 Torch 标杆复测。
135- 
136-| 场景 | 模型 | 特性 | Torch 小算子标杆 | 性能指标 | 测试环境 | 测试结果 | 遗留风险 |
137-| --- | --- | --- | --- | --- | --- | --- | --- |
138-| 算子微基准 | 不涉及 | ComplexMul / ComplexMulGrad | 同设备逐元素 `wr=ur*vr-ui*vi` 等 | 前向约 2.7x~2.8x,反向约 3.5x~5.5x | Atlas A2 + CANN 9.0.0 | Pass | 无 |
139-| 算子微基准 | 不涉及 | FusedBiasSoftmax | 同设备 `softmax(x+bias1+bias2,-1)`(fp32 累加) | bf16 约 2.7x~2.9x,小 fp32 `[2,4,64,64]` 约 1.66x | Atlas A2 + CANN 9.0.0 | Pass | 无 |
140-| 算子微基准 | 不涉及 | FusedSoftmaxGrad | 同设备 `g=go@V^T` + softmax JVP | bf16 约 1.98x~2.04x,fp32 约 1.14x | Atlas A2 + CANN 9.0.0 | Pass | 无 |
141-| 算子微基准 | 不涉及 | real_rfft / real_irfft | 同设备 `torch.fft.rfft` / `irfft`,`norm="ortho"` | 常规末轴约 1.1x~2.0x;业务 5D W/H/D 前向约 3.8x / 16.6x / 14.6x,irfft 约 4.1x / 18.2x / 14.5x;fwd+bwd 约 9.4x / 32x / 29x | Atlas A2 + CANN 9.0.0 | Pass | 无 |
142- 
143-代表形态实测如下(精度全部 PASS)。
144- 
145-**ComplexMul / ComplexMulGrad(相对同设备逐元素公式)**
146- 
147-| 用例 | 标杆 ms | 本算子 ms | 加速比 |
148-| --- | ---: | ---: | ---: |
149-| fwd `[4,1024]` | 0.270 | 0.102 | 2.65x |
150-| bwd `[4,1024]` | 0.652 | 0.119 | 5.49x |
151-| fwd `[1,1,32,32,32]` | 0.282 | 0.103 | 2.75x |
152-| bwd `[1,1,32,32,32]` | 0.644 | 0.126 | 5.13x |
153-| fwd `[1,1,96,160,160]` | 0.288 | 0.107 | 2.69x |
154-| bwd `[1,1,96,160,160]` | 0.652 | 0.126 | 5.17x |
155-| fwd `[1,1,192,320,320]` | 1.162 | 0.412 | 2.82x |
156-| bwd `[1,1,192,320,320]` | 2.478 | 0.709 | 3.50x |
157- 
158-**FusedBiasSoftmax(相对同设备 fp32 `softmax`)**
159- 
160-| 用例 | 标杆 ms | 本算子 ms | 加速比 |
161-| --- | ---: | ---: | ---: |
162-| fwd `[2,4,64,128]` bf16 | 0.529 | 0.183 | 2.89x |
163-| fwd `[2,4,256,256]` bf16 | 0.529 | 0.198 | 2.67x |
164-| fwd `[2,4,64,64]` bf16 | 0.240 | 0.085 | 2.83x |
165-| fwd `[2,4,64,64]` fp32 | 0.145 | 0.087 | 1.66x |
166-| fwd `[8,8,256,256]` bf16 | 0.244 | 0.091 | 2.68x |
167- 
168-**FusedSoftmaxGrad(相对同设备 `matmul` + softmax JVP)**
169- 
170-| 用例 | 标杆 ms | 本算子 ms | 加速比 |
171-| --- | ---: | ---: | ---: |
172-| bwd `[2,4,64,64]` D64 bf16 | 0.632 | 0.310 | 2.04x |
173-| bwd `[8,8,256,256]` D64 bf16 | 0.650 | 0.328 | 1.98x |
174-| bwd `[2,2,128,128]` D20 bf16 | 0.616 | 0.312 | 1.98x |
175-| bwd `[2,4,64,128]` D64 bf16 | 0.613 | 0.306 | 2.00x |
176-| bwd `[2,4,64,64]` D64 fp32 | 0.388 | 0.342 | 1.14x |
177- 
178-**real_rfft / real_irfft 常规形态(相对同设备 `torch.fft`)**
179- 
180-相对误差约 `2e-7`~`5e-7`。含业务 5D 的全表几何平均加速比约 **2.71x**。业务 5D 相对标杆少约 1.1~1.7 GB 峰值显存。
181- 
182-| 用例 | 标杆 ms | 本算子 ms | 加速比 |
183-| --- | ---: | ---: | ---: |
184-| rfft `[8,256]` 末轴 | 0.401 | 0.275 | 1.46x |
185-| irfft `[8,256]` 末轴 | 0.554 | 0.280 | 1.98x |
186-| rfft `[32,512]` 末轴 | 0.406 | 0.260 | 1.56x |
187-| irfft `[32,512]` 末轴 | 0.546 | 0.272 | 2.01x |
188-| rfft `128³` 末轴 | 0.300 | 0.262 | 1.14x |
189-| irfft `128³` 末轴 | 0.488 | 0.283 | 1.73x |
190-| rfft `128³` 非末轴 | 0.270 | 0.271 | 1.00x |
191-| irfft `128³` 非末轴 | 0.478 | 0.284 | 1.68x |
192- 
193-**real_rfft / real_irfft 业务 5D `[1,35,128,128,128]`(相对同设备 `torch.fft`)**
194- 
195-| 轴 | rfft | irfft | fwd+bwd |
196-| --- | ---: | ---: | ---: |
197-| W 末轴 | 3.81x(7.964 / 2.089 ms) | 4.11x(10.573 / 2.570 ms) | 9.39x(45.875 / 4.886 ms) |
198-| H | 16.59x(8.501 / 0.512 ms) | 18.20x(10.517 / 0.578 ms) | 32.23x(49.057 / 1.522 ms) |
199-| D | 14.59x(7.908 / 0.542 ms) | 14.53x(10.646 / 0.733 ms) | 29.15x(48.925 / 1.678 ms) |
200- 
201-### 5.4 兼容性测试
202- 
203-兼容性评估:通过。
204- 
205-| 序号 | 兼容性场景 | 验证结果 | 遗留风险 |
206-| --- | --- | --- | --- |
207-| 1 | CANN 9.0.0 配套(Toolkit + 910B Ops) | pass | 无 |
208-| 2 | Python 3.11 + PyTorch 2.7.1 + 配套 torch_npu | pass | 无 |
209-| 3 | `requirements.txt` 声明的构建与测试依赖 | pass | 无 |
210-| 4 | docker 与 host+device 形态下编译、安装、跑通算子 | pass | 无 |
211-| 5 | 文档声明的 Atlas A2 / Ascend 910B | pass | Atlas A3 / 950 系列文档标明暂未验证 |
212- 
213-## 6. 测试执行评估
214- 
215-### 6.1 测试覆盖
216- 
217-特性以及关键测试活动的测试覆盖情况如下。
218- 
219-| 测试活动 | 测试结论 | 用例数 | 用例覆盖率 | 用例通过率 |
220-| --- | --- | --- | --- | --- |
221-| 特性测试 | pass | 16 | 100% | 100% |
222-| 性能测试 | pass | 4 | 100% | 100% |
223-| 可靠性测试 | pass | 4 | 100% | 100% |
224-| 继承特性测试 | 不涉及 | 0 | — | — |
225-| 兼容性测试 | pass | 5 | 100% | 100% |
226-| 安全测试 | pass | 2 | 100% | 100% |
227- 
228-说明:特性测试第 1~16 组覆盖文档、编译安装与全部已交付算子;性能测试、可靠性、兼容性、安全与第 4、5 节表格对应。首发版本无继承特性。下游模型仓交叉验证未配置,不纳入本次范围,不计入失败。
229- 
230-## 7. 遗留问题和关键风险
231- 
232-不涉及本仓功能缺陷。
233- 
234-### 7.1 遗留问题统计
235- 
236-| | 问题总数 | 严重 | 主要 | 次要 | 不重要 | 已取消 |
237-| --- | --- | --- | --- | --- | --- | --- |
238-| 数目 | 0 | 0 | 0 | 0 | 0 | 0 |
239-| 百分比 | 100% | 0% | 0% | 0% | 0% | 0% |
240- 
241-### 7.2 遗留问题列表
242- 
243-| 问题单(issue链接) | 问题描述 | 问题级别 | 问题影响和规避措施 | 当前状态 |
244-| --- | --- | --- | --- | --- |
245-| 不涉及 | 不涉及 | — | — | — |
246- 
247-已知环境差异(非本仓库失败,不计入上表):原生 `torch.fft` 在部分 NPU 官方实现上的反向与 CPU 不完全一致;本仓 `real_rfft` / `real_irfft` 以 CPU `torch.fft` 为标杆验收。官方用例门槛 `1e-5`;相位折叠造基后实测相对误差约 `2e-7`~`5e-7`。
248- 
249-## 8. 附件
250-不涉及
@@ -53,7 +53,7 @@ w_i &= u_r v_i + u_i v_r
53- 四个输入与两个输出必须 **同 shape、同 dtype、`float32`、`ND`**。53- 四个输入与两个输出必须 **同 shape、同 dtype、`float32`、`ND`**。
54- 输入需连续(Python 绑定会 `contiguous()`)。54- 输入需连续(Python 绑定会 `contiguous()`)。
55- **AIV-only**:按元素总数在 AIV 核间均分(32B 对齐),核内再按 UB 预算分拍。前向 UB 模型为 13 份 tile(4 路 inQue + 2 路 outQue 的 double buffer + 1 份 tmp)。55- **AIV-only**:按元素总数在 AIV 核间均分(32B 对齐),核内再按 UB 预算分拍。前向 UB 模型为 13 份 tile(4 路 inQue + 2 路 outQue 的 double buffer + 1 份 tmp)。
56-- 切分算法在 `op_host/complex_mul_tiling_core.h`,由 [`tests/ascendc/ophost/test_complex_mul_tiling.cpp`](../../tests/ascendc/ophost/test_complex_mul_tiling.cpp) 覆盖。56+- 切分算法在 `op_host/complex_mul_tiling_core.h`,由 [`tests/ut/op_host/complex_mul/test_complex_mul_tiling.cpp`](../../tests/ut/op_host/complex_mul/test_complex_mul_tiling.cpp) 覆盖。
57- Python 默认走 PyTorch 参考实现;仅当张量在 NPU 且 `OIL_GAS_OPS_COMPLEX_MUL_IMPL` 为 `ascendc` / `auto` 时调用本算子。57- Python 默认走 PyTorch 参考实现;仅当张量在 NPU 且 `OIL_GAS_OPS_COMPLEX_MUL_IMPL` 为 `ascendc` / `auto` 时调用本算子。
58- 精度验收:相对误差默认阈值 `1e-5`(见 [`tests/ascendc/complex_mul/`](../../tests/ascendc/complex_mul/))。58- 精度验收:相对误差默认阈值 `1e-5`(见 [`tests/ascendc/complex_mul/`](../../tests/ascendc/complex_mul/))。
59 59 
@@ -126,7 +126,7 @@ NPU 融合路径上 **`y` 与 `x` 共享同一块 device 内存**(原地写回
126- 存在的 bias 必须与 `x` **同 dtype**。126- 存在的 bias 必须与 `x` **同 dtype**。
127- **AIV-only**。按行(`B*H*Q`)在 AIV 核间均分,每行 `S` 对齐到 32B;高阶 SoftMax 预留 32KB UB。127- **AIV-only**。按行(`B*H*Q`)在 AIV 核间均分,每行 `S` 对齐到 32B;高阶 SoftMax 预留 32KB UB。
128- tiling key:`0=fp32`,`1=fp16`,`2=bf16`。128- tiling key:`0=fp32`,`1=fp16`,`2=bf16`。
129-- 切分算法在 `op_host/fused_bias_softmax_tiling_core.h`,由 [`tests/ascendc/ophost/test_fused_bias_softmax_tiling.cpp`](../../tests/ascendc/ophost/test_fused_bias_softmax_tiling.cpp) 覆盖。129+- 切分算法在 `op_host/fused_bias_softmax_tiling_core.h`,由 [`tests/ut/op_host/fused_bias_softmax/test_fused_bias_softmax_tiling.cpp`](../../tests/ut/op_host/fused_bias_softmax/test_fused_bias_softmax_tiling.cpp) 覆盖。
130- 精度验收([`tests/ascendc/fused_bias_softmax/`](../../tests/ascendc/fused_bias_softmax/)):fp32 `max_diff ≤ 1e-5`;bf16 契约式 `atol=2e-2, rtol=4e-3`。130- 精度验收([`tests/ascendc/fused_bias_softmax/`](../../tests/ascendc/fused_bias_softmax/)):fp32 `max_diff ≤ 1e-5`;bf16 契约式 `atol=2e-2, rtol=4e-3`。
131 131 
132---132---
@@ -175,7 +175,7 @@ NPU 融合路径上 **`out` 与 `probs` 共享同一块 device 内存**(原地
175- `fp32` / `fp16` 不走 cube:Ascend 910B fixpipe C 为 16bit 结构上限,`g` 以 half 落盘的量化误差会超出 fp32 契约(`rtol=1e-4`)。175- `fp32` / `fp16` 不走 cube:Ascend 910B fixpipe C 为 16bit 结构上限,`g` 以 half 落盘的量化误差会超出 fp32 契约(`rtol=1e-4`)。
176- 向量路径要求 \(V\) 的 plane(`S * pad(D) * 4` 字节)加上约 48KB 余量能放入 AIV UB,否则 tiling 失败。176- 向量路径要求 \(V\) 的 plane(`S * pad(D) * 4` 字节)加上约 48KB 余量能放入 AIV UB,否则 tiling 失败。
177- MIX 路径 workspace = 系统预留 16MB + `B*H*Q*S*sizeof(half)` 的 `g` 缓冲。177- MIX 路径 workspace = 系统预留 16MB + `B*H*Q*S*sizeof(half)` 的 `g` 缓冲。
178-- 切分算法在 `op_host/fused_softmax_grad_tiling_core.h`,由 [`tests/ascendc/ophost/test_fused_softmax_grad_tiling.cpp`](../../tests/ascendc/ophost/test_fused_softmax_grad_tiling.cpp) 覆盖。178+- 切分算法在 `op_host/fused_softmax_grad_tiling_core.h`,由 [`tests/ut/op_host/fused_softmax_grad/test_fused_softmax_grad_tiling.cpp`](../../tests/ut/op_host/fused_softmax_grad/test_fused_softmax_grad_tiling.cpp) 覆盖。
179- 精度验收与前向同一目录:fp32 `max_diff ≤ 1e-5`;bf16 `atol=2e-2, rtol=4e-3`。MIX 路径的 half-`g` 量化误差经 `rowsum(P⊙g)` 放大,与 \(\sum P^2\) 成正比,生产输入应为 softmax 分布。179- 精度验收与前向同一目录:fp32 `max_diff ≤ 1e-5`;bf16 `atol=2e-2, rtol=4e-3`。MIX 路径的 half-`g` 量化误差经 `rowsum(P⊙g)` 放大,与 \(\sum P^2\) 成正比,生产输入应为 softmax 分布。
180 180 
181`attention_core(q, k, v, bias1, bias2)` 是上述两算子按 openfold `AttentionCoreFunction.apply` 语义的组合封装,不是独立 OPP 算子。数学定义见[接口列表](api_list.md#attention_core)。181`attention_core(q, k, v, bias1, bias2)` 是上述两算子按 openfold `AttentionCoreFunction.apply` 语义的组合封装,不是独立 OPP 算子。数学定义见[接口列表](api_list.md#attention_core)。
@@ -4,12 +4,20 @@
4 4 
5| 示例 | 说明 | 需要 NPU |5| 示例 | 说明 | 需要 NPU |
6|------|------|----------|6|------|------|----------|
7-| [`complex_mul_example.py`](complex_mul_example.py) | AscendC 复数逐点乘的前向、反向与精度对拍 | 否(无卡时走 PyTorch 参考实现) |7+| [`complex_mul_example.py`](complex_mul_example.py) | AscendC 四路 float 复数逐点乘的前向、反向与精度对拍 | 有 NPU 时走自定义算子,否则走 PyTorch 参考实现 |
8-| [`real_fft_example.py`](real_fft_example.py) | Triton-Ascend 实数 FFT 的正反变换、往返重建与 autograd | 否(无卡时回退到 `torch.fft`) |8+| [`real_fft_example.py`](real_fft_example.py) | Triton-Ascend 实数 FFT 的正反变换、往返重建与 autograd | 有 NPU 时走 Triton kernel,否则回退 `torch.fft` |
9+ 
10+各 AscendC 算子还有对齐 ops-nn addmv 的 C++ 两段式样例:`examples/aclnn/<op>/test_aclnn_*.cpp`(需要 NPU + customize vendor)。
11+ 
12+```bash
13+bash build.sh -u --examples
14+# 或单算子:
15+bash examples/aclnn/complex_mul/run.sh
16+```
9 17 
10## 运行18## 运行
11 19 
12-两个示例都会自动探测设备:有 NPU 时走自定义算子,无 NPU 时走各自的回退实现,因此在开发容器里也能跑通。未安装 whl 时会自动回落到仓内源码(见 [`_bootstrap.py`](_bootstrap.py)),不需要额外设置 `PYTHONPATH`:20+两个示例都会自动探测设备:有 NPU 时走自定义算子,否则走各自的回退实现。未安装 whl 时会自动回落到仓内源码(见 [`_bootstrap.py`](_bootstrap.py)),不需要额外设置 `PYTHONPATH`:
13 21 
14```bash22```bash
15python3 examples/complex_mul_example.py23python3 examples/complex_mul_example.py
@@ -0,0 +1,76 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# Licensed under the Apache License, Version 2.0 (the "License");
4+# you may not use this file except in compliance with the License.
5+# You may obtain a copy of the License at
6+#
7+# http://www.apache.org/licenses/LICENSE-2.0
8+#
9+# Unless required by applicable law or agreed to in writing, software
10+# distributed under the License is distributed on an "AS IS" BASIS,
11+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+# See the License for the specific language governing permissions and
13+# limitations under the License.
14+# -----------------------------------------------------------------------------------------------------------
15+#
16+# 对齐 ops-nn <op>/examples/test_aclnn_*.cpp 的独立编译。
17+# 调用方须先 project() 并设置 ACLNN_TARGET、ACLNN_EXAMPLE_SRC。
18+ 
19+if(NOT ACLNN_TARGET)
20+ message(FATAL_ERROR "ACLNN_TARGET must be set before including aclnn_example.cmake")
21+endif()
22+if(NOT ACLNN_EXAMPLE_SRC)
23+ message(FATAL_ERROR "ACLNN_EXAMPLE_SRC must be set before including aclnn_example.cmake")
24+endif()
25+ 
26+add_compile_options(-std=c++17)
27+set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/output")
28+ 
29+set(INC_PATH "$ENV{DDK_PATH}")
30+if(INC_PATH STREQUAL "")
31+ if(DEFINED ENV{ASCEND_HOME_PATH} AND NOT "$ENV{ASCEND_HOME_PATH}" STREQUAL "")
32+ set(INC_PATH "$ENV{ASCEND_HOME_PATH}")
33+ else()
34+ set(INC_PATH "/usr/local/Ascend/ascend-toolkit/latest")
35+ endif()
36+ message(STATUS "set default INC_PATH: ${INC_PATH}")
37+else()
38+ message(STATUS "env INC_PATH: ${INC_PATH}")
39+endif()
40+ 
41+set(CUST_PKG_PATH "${INC_PATH}/opp/vendors/customize/op_api")
42+set(LIB_PATH "$ENV{NPU_HOST_LIB}")
43+if(LIB_PATH STREQUAL "")
44+ string(TOLOWER "${CMAKE_SYSTEM_NAME}" SYSTEM_NAME_LOWER)
45+ set(LIB_PATH "${INC_PATH}/${CMAKE_SYSTEM_PROCESSOR}-${SYSTEM_NAME_LOWER}/lib64")
46+ message(STATUS "set default LIB_PATH: ${LIB_PATH}")
47+else()
48+ message(STATUS "env LIB_PATH: ${LIB_PATH}")
49+endif()
50+ 
51+get_filename_component(ACLNN_EXAMPLE_DIR "${CMAKE_CURRENT_LIST_DIR}" ABSOLUTE)
52+get_filename_component(OG_REPO_ROOT "${ACLNN_EXAMPLE_DIR}/../.." ABSOLUTE)
53+ 
54+include_directories(
55+ ${ACLNN_EXAMPLE_DIR}
56+ ${OG_REPO_ROOT}/ascendc/operators/_common/op_api
57+ ${INC_PATH}/include
58+ ${CUST_PKG_PATH}/include
59+)
60+ 
61+link_directories(
62+ ${LIB_PATH}
63+ ${CUST_PKG_PATH}/lib
64+)
65+ 
66+add_executable(${ACLNN_TARGET} ${ACLNN_EXAMPLE_SRC})
67+ 
68+target_link_libraries(${ACLNN_TARGET}
69+ ascendcl
70+ cust_opapi
71+ acl_op_compiler
72+ nnopbase
73+ stdc++
74+)
75+ 
76+install(TARGETS ${ACLNN_TARGET} DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY})
@@ -0,0 +1,234 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * Licensed under the Apache License, Version 2.0 (the "License");
4+ * you may not use this file except in compliance with the License.
5+ * You may obtain a copy of the License at
6+ *
7+ * http://www.apache.org/licenses/LICENSE-2.0
8+ *
9+ * Unless required by applicable law or agreed to in writing, software
10+ * distributed under the License is distributed on an "AS IS" BASIS,
11+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+ * See the License for the specific language governing permissions and
13+ * limitations under the License.
14+ */
15+ 
16+/**
17+ * 对齐 ops-nn matmul/addmv/examples/test_aclnn_addmv.cpp 的两段式 aclnn 小样例脚手架。
18+ * 各算子 examples/test_aclnn_*.cpp 内联构造小 tensor,host 侧核对结果。
19+ */
20+ 
21+#ifndef OG_ACLNN_EXAMPLE_H
22+#define OG_ACLNN_EXAMPLE_H
23+ 
24+#include <cmath>
25+#include <cstdint>
26+#include <cstdio>
27+#include <iostream>
28+#include <vector>
29+ 
30+#include "acl/acl.h"
31+#include "aclnn/acl_meta.h"
32+ 
33+#define CHECK_RET(cond, return_expr) \
34+ do { \
35+ if (!(cond)) { \
36+ return_expr; \
37+ } \
38+ } while (0)
39+ 
40+#define LOG_PRINT(message, ...) \
41+ do { \
42+ printf(message, ##__VA_ARGS__); \
43+ } while (0)
44+ 
45+inline int64_t GetShapeSize(const std::vector<int64_t>& shape)
46+{
47+ int64_t shapeSize = 1;
48+ for (auto dim : shape) {
49+ shapeSize *= dim;
50+ }
51+ return shapeSize;
52+}
53+ 
54+inline std::vector<int64_t> ContiguousStrides(const std::vector<int64_t>& shape)
55+{
56+ std::vector<int64_t> strides(shape.size(), 1);
57+ for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) {
58+ strides[static_cast<size_t>(i)] = shape[static_cast<size_t>(i + 1)] * strides[static_cast<size_t>(i + 1)];
59+ }
60+ return strides;
61+}
62+ 
63+inline int Init(int32_t deviceId, aclrtStream* stream)
64+{
65+ auto ret = aclInit(nullptr);
66+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
67+ ret = aclrtSetDevice(deviceId);
68+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
69+ ret = aclrtCreateStream(stream);
70+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
71+ return 0;
72+}
73+ 
74+inline aclTensor* MakeNdTensor(const std::vector<int64_t>& shape, aclDataType dataType, void* deviceAddr,
75+ int64_t storageNumel = -1)
76+{
77+ const auto strides = ContiguousStrides(shape);
78+ const int64_t numel = storageNumel > 0 ? storageNumel : GetShapeSize(shape);
79+ std::vector<int64_t> storageDims = {numel};
80+ return aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
81+ storageDims.data(), 1, deviceAddr);
82+}
83+ 
84+template <typename T>
85+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
86+ aclDataType dataType, aclTensor** tensor, int64_t storageNumel = -1)
87+{
88+ const int64_t logical = GetShapeSize(shape);
89+ const int64_t storage = storageNumel > 0 ? storageNumel : static_cast<int64_t>(hostData.size());
90+ CHECK_RET(storage >= logical, LOG_PRINT("hostData smaller than logical shape\n"); return 1);
91+ const size_t bytes = static_cast<size_t>(storage) * sizeof(T);
92+ auto ret = aclrtMalloc(deviceAddr, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
93+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
94+ ret = aclrtMemcpy(*deviceAddr, bytes, hostData.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE);
95+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
96+ *tensor = MakeNdTensor(shape, dataType, *deviceAddr, storage);
97+ return 0;
98+}
99+ 
100+inline int CreateAclTensorView(void* deviceAddr, const std::vector<int64_t>& shape, aclDataType dataType,
101+ aclTensor** tensor, int64_t storageNumel = -1)
102+{
103+ *tensor = MakeNdTensor(shape, dataType, deviceAddr, storageNumel);
104+ return *tensor == nullptr ? 1 : 0;
105+}
106+ 
107+inline bool HostClose(const std::vector<float>& actual, const std::vector<float>& expected, float atol = 1e-5f,
108+ float rtol = 1e-5f)
109+{
110+ if (actual.size() != expected.size()) {
111+ LOG_PRINT("size mismatch: actual=%zu expected=%zu\n", actual.size(), expected.size());
112+ return false;
113+ }
114+ for (size_t i = 0; i < actual.size(); ++i) {
115+ const float exp = expected[i];
116+ const float got = actual[i];
117+ const float tol = atol + rtol * std::fabs(exp);
118+ if (std::fabs(got - exp) > tol) {
119+ LOG_PRINT("mismatch[%zu] expected=%f actual=%f\n", i, exp, got);
120+ return false;
121+ }
122+ }
123+ return true;
124+}
125+ 
126+inline int CopyOut(void* deviceAddr, std::vector<float>& host, size_t count)
127+{
128+ host.resize(count);
129+ const size_t bytes = count * sizeof(float);
130+ auto ret = aclrtMemcpy(host.data(), bytes, deviceAddr, bytes, ACL_MEMCPY_DEVICE_TO_HOST);
131+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
132+ return 0;
133+}
134+ 
135+/** workspaceSize 为 0 时不分配,*workspaceAddr 保持 nullptr。 */
136+inline int MallocWorkspace(uint64_t workspaceSize, void** workspaceAddr)
137+{
138+ if (workspaceSize == 0) {
139+ return ACL_SUCCESS;
140+ }
141+ auto ret = aclrtMalloc(workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
142+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
143+ return ACL_SUCCESS;
144+}
145+ 
146+inline void PrintResult(const std::vector<float>& result)
147+{
148+ for (size_t i = 0; i < result.size(); ++i) {
149+ LOG_PRINT("result[%zu] is: %f\n", i, result[i]);
150+ }
151+}
152+ 
153+inline void DestroyTensors(aclTensor** tensors, size_t count)
154+{
155+ for (size_t i = 0; i < count; ++i) {
156+ if (tensors[i] != nullptr) {
157+ (void)aclDestroyTensor(tensors[i]);
158+ tensors[i] = nullptr;
159+ }
160+ }
161+}
162+ 
163+inline void FreeDeviceBuffers(void** buffers, size_t count)
164+{
165+ for (size_t i = 0; i < count; ++i) {
166+ if (buffers[i] != nullptr) {
167+ (void)aclrtFree(buffers[i]);
168+ buffers[i] = nullptr;
169+ }
170+ }
171+}
172+ 
173+/** 释放 count 个输入 device 缓冲,外加一块打包输出缓冲。 */
174+inline void FreePackedBuffers(void** inputs, size_t count, void** packed)
175+{
176+ FreeDeviceBuffers(inputs, count);
177+ FreeDeviceBuffers(packed, 1);
178+}
179+ 
180+/** 样例收尾:释放 workspace、销毁 stream、复位 device、去初始化 acl。 */
181+inline void Finalize(int32_t deviceId, aclrtStream stream, void* workspaceAddr)
182+{
183+ if (workspaceAddr != nullptr) {
184+ (void)aclrtFree(workspaceAddr);
185+ }
186+ (void)aclrtDestroyStream(stream);
187+ (void)aclrtResetDevice(deviceId);
188+ (void)aclFinalize();
189+}
190+ 
191+/**
192+ * 各样例的算子相关回调。TensorsT 由样例自定义,持有 host 输入、device 缓冲与 aclTensor。
193+ * setup 建 tensor,launch 下发两段式接口,outputAddr 指向回读用的 device 缓冲成员,
194+ * expected 给出 host 侧期望值,release 回收资源。
195+ */
196+template <typename TensorsT>
197+struct AclnnExampleOps {
198+ int (*setup)(TensorsT&);
199+ int (*launch)(const TensorsT&, aclrtStream, void**);
200+ void* TensorsT::*outputAddr;
201+ std::vector<float> (*expected)(const TensorsT&);
202+ void (*release)(TensorsT&);
203+};
204+ 
205+/** 样例主流程:Init -> setup -> launch -> 回读 -> 对拍 -> 打印 -> 释放 -> 收尾。对拍不过返回 1。 */
206+template <typename TensorsT>
207+int RunAclnnExample(const AclnnExampleOps<TensorsT>& ops, size_t outNumel, float atol = 1e-5f, float rtol = 1e-5f)
208+{
209+ int32_t deviceId = 0;
210+ aclrtStream stream = nullptr;
211+ auto ret = Init(deviceId, &stream);
212+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
213+ 
214+ TensorsT tensors;
215+ void* workspaceAddr = nullptr;
216+ ret = ops.setup(tensors);
217+ CHECK_RET(ret == ACL_SUCCESS, return ret);
218+ ret = ops.launch(tensors, stream, &workspaceAddr);
219+ CHECK_RET(ret == ACL_SUCCESS, return ret);
220+ 
221+ std::vector<float> result;
222+ ret = CopyOut(tensors.*ops.outputAddr, result, outNumel);
223+ CHECK_RET(ret == ACL_SUCCESS, return ret);
224+ 
225+ const bool close = HostClose(result, ops.expected(tensors), atol, rtol);
226+ if (close) {
227+ PrintResult(result);
228+ }
229+ ops.release(tensors);
230+ Finalize(deviceId, stream, workspaceAddr);
231+ return close ? 0 : 1;
232+}
233+ 
234+#endif
@@ -0,0 +1,56 @@
1+#!/bin/bash
2+# -----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# Licensed under the Apache License, Version 2.0 (the "License");
5+# you may not use this file except in compliance with the License.
6+# You may obtain a copy of the License at
7+#
8+# http://www.apache.org/licenses/LICENSE-2.0
9+#
10+# Unless required by applicable law or agreed to in writing, software
11+# distributed under the License is distributed on an "AS IS" BASIS,
12+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+# See the License for the specific language governing permissions and
14+# limitations under the License.
15+# -----------------------------------------------------------------------------------------------------------
16+#
17+# 由 examples/aclnn/<op>/run.sh source,需已设置:
18+# CURRENT_DIR 该算子 examples/aclnn/<op> 目录
19+# ACLNN_EXE 可执行文件名(与 CMake ACLNN_TARGET 一致)
20+ 
21+set -u
22+ 
23+EXAMPLE_COMMON_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
24+REPO_ROOT="$(cd "${EXAMPLE_COMMON_DIR}/../.." && pwd)"
25+ 
26+if [ -z "${CURRENT_DIR:-}" ] || [ -z "${ACLNN_EXE:-}" ]; then
27+ echo "[ERROR]: CURRENT_DIR and ACLNN_EXE must be set" >&2
28+ exit 1
29+fi
30+ 
31+# shellcheck source-path=SCRIPTDIR source=../../scripts/cann_env.sh
32+source "${REPO_ROOT}/scripts/cann_env.sh"
33+cann_setup_env
34+ 
35+example_stage() {
36+ local stage="$1"
37+ shift
38+ if ! "$@"; then
39+ echo "[ERROR]: example stage '${stage}' failed for ${ACLNN_EXE}" >&2
40+ exit 1
41+ fi
42+}
43+ 
44+rm -rf "${CURRENT_DIR}/build" "${CURRENT_DIR}/output"
45+mkdir -p "${CURRENT_DIR}/build"
46+ 
47+cd "${CURRENT_DIR}/build" || exit 1
48+example_stage cmake cmake .. -DCMAKE_SKIP_RPATH=TRUE
49+example_stage make make
50+ 
51+cann_export_custom_op_api
52+ 
53+cd "${CURRENT_DIR}/output" || exit 1
54+echo "[INFO]: Execute example ${ACLNN_EXE}"
55+example_stage launch "./${ACLNN_EXE}"
56+echo "[INFO]: example success"
@@ -0,0 +1,20 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# Licensed under the Apache License, Version 2.0 (the "License");
4+# you may not use this file except in compliance with the License.
5+# You may obtain a copy of the License at
6+#
7+# http://www.apache.org/licenses/LICENSE-2.0
8+#
9+# Unless required by applicable law or agreed to in writing, software
10+# distributed under the License is distributed on an "AS IS" BASIS,
11+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+# See the License for the specific language governing permissions and
13+# limitations under the License.
14+# -----------------------------------------------------------------------------------------------------------
15+ 
16+cmake_minimum_required(VERSION 3.16)
17+project(test_aclnn_complex_mul)
18+set(ACLNN_TARGET test_aclnn_complex_mul)
19+set(ACLNN_EXAMPLE_SRC test_aclnn_complex_mul.cpp)
20+include("${CMAKE_CURRENT_SOURCE_DIR}/../../_common/aclnn_example.cmake")
@@ -0,0 +1,23 @@
1+#!/bin/bash
2+# -----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# Licensed under the Apache License, Version 2.0 (the "License");
5+# you may not use this file except in compliance with the License.
6+# You may obtain a copy of the License at
7+#
8+# http://www.apache.org/licenses/LICENSE-2.0
9+#
10+# Unless required by applicable law or agreed to in writing, software
11+# distributed under the License is distributed on an "AS IS" BASIS,
12+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+# See the License for the specific language governing permissions and
14+# limitations under the License.
15+# -----------------------------------------------------------------------------------------------------------
16+ 
17+CURRENT_DIR=$(
18+ cd $(dirname ${BASH_SOURCE:-$0})
19+ pwd
20+)
21+ACLNN_EXE=test_aclnn_complex_mul
22+# shellcheck source=../../_common/run_example.sh
23+source "${CURRENT_DIR}/../../_common/run_example.sh"
@@ -0,0 +1,97 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * Licensed under the Apache License, Version 2.0 (the "License");
4+ * you may not use this file except in compliance with the License.
5+ * You may obtain a copy of the License at
6+ *
7+ * http://www.apache.org/licenses/LICENSE-2.0
8+ *
9+ * Unless required by applicable law or agreed to in writing, software
10+ * distributed under the License is distributed on an "AS IS" BASIS,
11+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+ * See the License for the specific language governing permissions and
13+ * limitations under the License.
14+ */
15+ 
16+/**
17+ * 对齐 ops-nn addmv/examples:内联小 tensor,两段式 aclnnComplexMul。
18+ * wr/wi 共用一块 packed 缓冲(kernel 忽略独立 wi 指针)。
19+ */
20+ 
21+#include "aclnn_example.h"
22+#include "og_aclnn_decls.h"
23+ 
24+namespace {
25+ 
26+constexpr size_t kNumel = 4;
27+constexpr size_t kPackedNumel = kNumel * 2;
28+constexpr size_t kInputCount = 4;
29+ 
30+const std::vector<int64_t> kShape = {2, 2};
31+const std::vector<float> kUr = {1, 2, 3, 4};
32+const std::vector<float> kUi = {0, 1, -1, 0.5f};
33+const std::vector<float> kVr = {0.5f, 1, 1, 2};
34+const std::vector<float> kVi = {1, 0, 0.5f, -1};
35+ 
36+struct Tensors {
37+ void* inDev[kInputCount] = {};
38+ aclTensor* inT[kInputCount] = {};
39+ void* packedDev = nullptr;
40+ aclTensor* wrT = nullptr;
41+ aclTensor* wiT = nullptr;
42+};
43+ 
44+int SetupTensors(Tensors& t)
45+{
46+ const std::vector<float>* hosts[kInputCount] = {&kUr, &kUi, &kVr, &kVi};
47+ for (size_t i = 0; i < kInputCount; ++i) {
48+ auto ret = CreateAclTensor(*hosts[i], kShape, &t.inDev[i], aclDataType::ACL_FLOAT, &t.inT[i]);
49+ CHECK_RET(ret == ACL_SUCCESS, return ret);
50+ }
51+ const std::vector<float> packed(kPackedNumel, 0.0f);
52+ auto ret = CreateAclTensor(packed, kShape, &t.packedDev, aclDataType::ACL_FLOAT, &t.wrT, kPackedNumel);
53+ CHECK_RET(ret == ACL_SUCCESS, return ret);
54+ return CreateAclTensorView(t.packedDev, kShape, aclDataType::ACL_FLOAT, &t.wiT);
55+}
56+ 
57+int LaunchOp(const Tensors& t, aclrtStream stream, void** workspaceAddr)
58+{
59+ uint64_t workspaceSize = 0;
60+ aclOpExecutor* executor = nullptr;
61+ auto ret = aclnnComplexMulGetWorkspaceSize(t.inT[0], t.inT[1], t.inT[2], t.inT[3], t.wrT, t.wiT, &workspaceSize,
62+ &executor);
63+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnComplexMulGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
64+ ret = MallocWorkspace(workspaceSize, workspaceAddr);
65+ CHECK_RET(ret == ACL_SUCCESS, return ret);
66+ ret = aclnnComplexMul(*workspaceAddr, workspaceSize, executor, stream);
67+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnComplexMul failed. ERROR: %d\n", ret); return ret);
68+ ret = aclrtSynchronizeStream(stream);
69+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
70+ return ACL_SUCCESS;
71+}
72+ 
73+std::vector<float> Expected(const Tensors&)
74+{
75+ std::vector<float> expected(kPackedNumel, 0.0f);
76+ for (size_t i = 0; i < kNumel; ++i) {
77+ expected[i] = kUr[i] * kVr[i] - kUi[i] * kVi[i];
78+ expected[i + kNumel] = kUr[i] * kVi[i] + kUi[i] * kVr[i];
79+ }
80+ return expected;
81+}
82+ 
83+void Release(Tensors& t)
84+{
85+ DestroyTensors(t.inT, kInputCount);
86+ DestroyTensors(&t.wrT, 1);
87+ DestroyTensors(&t.wiT, 1);
88+ FreePackedBuffers(t.inDev, kInputCount, &t.packedDev);
89+}
90+ 
91+} // namespace
92+ 
93+int main()
94+{
95+ const AclnnExampleOps<Tensors> ops{SetupTensors, LaunchOp, &Tensors::packedDev, Expected, Release};
96+ return RunAclnnExample(ops, kPackedNumel);
97+}
@@ -0,0 +1,20 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# Licensed under the Apache License, Version 2.0 (the "License");
4+# you may not use this file except in compliance with the License.
5+# You may obtain a copy of the License at
6+#
7+# http://www.apache.org/licenses/LICENSE-2.0
8+#
9+# Unless required by applicable law or agreed to in writing, software
10+# distributed under the License is distributed on an "AS IS" BASIS,
11+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+# See the License for the specific language governing permissions and
13+# limitations under the License.
14+# -----------------------------------------------------------------------------------------------------------
15+ 
16+cmake_minimum_required(VERSION 3.16)
17+project(test_aclnn_complex_mul_grad)
18+set(ACLNN_TARGET test_aclnn_complex_mul_grad)
19+set(ACLNN_EXAMPLE_SRC test_aclnn_complex_mul_grad.cpp)
20+include("${CMAKE_CURRENT_SOURCE_DIR}/../../_common/aclnn_example.cmake")
@@ -0,0 +1,23 @@
1+#!/bin/bash
2+# -----------------------------------------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# Licensed under the Apache License, Version 2.0 (the "License");
5+# you may not use this file except in compliance with the License.
6+# You may obtain a copy of the License at
7+#
8+# http://www.apache.org/licenses/LICENSE-2.0
9+#
10+# Unless required by applicable law or agreed to in writing, software
11+# distributed under the License is distributed on an "AS IS" BASIS,
12+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+# See the License for the specific language governing permissions and
14+# limitations under the License.
15+# -----------------------------------------------------------------------------------------------------------
16+ 
17+CURRENT_DIR=$(
18+ cd $(dirname ${BASH_SOURCE:-$0})
19+ pwd
20+)
21+ACLNN_EXE=test_aclnn_complex_mul_grad
22+# shellcheck source=../../_common/run_example.sh
23+source "${CURRENT_DIR}/../../_common/run_example.sh"
@@ -0,0 +1,106 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * Licensed under the Apache License, Version 2.0 (the "License");
4+ * you may not use this file except in compliance with the License.
5+ * You may obtain a copy of the License at
6+ *
7+ * http://www.apache.org/licenses/LICENSE-2.0
8+ *
9+ * Unless required by applicable law or agreed to in writing, software
10+ * distributed under the License is distributed on an "AS IS" BASIS,
11+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+ * See the License for the specific language governing permissions and
13+ * limitations under the License.
14+ */
15+ 
16+/**
17+ * 两段式 aclnnComplexMulGrad 小样例。
18+ * gur/gui/gvr/gvi 共用一块 packed 缓冲(kernel 忽略后三个独立指针)。
19+ */
20+ 
21+#include "aclnn_example.h"
22+#include "og_aclnn_decls.h"
23+ 
24+namespace {
25+ 
26+constexpr size_t kNumel = 4;
27+constexpr size_t kInputCount = 6;
28+constexpr size_t kOutputCount = 4;
29+constexpr size_t kPackedNumel = kNumel * kOutputCount;
30+ 
31+const std::vector<int64_t> kShape = {2, 2};
32+const std::vector<float> kGwr = {1, 0, 0.5f, -1};
33+const std::vector<float> kGwi = {0, 1, -0.5f, 0.25f};
34+const std::vector<float> kUr = {1, 2, 3, 4};
35+const std::vector<float> kUi = {0, 1, -1, 0.5f};
36+const std::vector<float> kVr = {0.5f, 1, 1, 2};
37+const std::vector<float> kVi = {1, 0, 0.5f, -1};
38+ 
39+struct Tensors {
40+ void* inDev[kInputCount] = {};
41+ aclTensor* inT[kInputCount] = {};
42+ void* packedDev = nullptr;
43+ aclTensor* outT[kOutputCount] = {};
44+};
45+ 
46+int SetupTensors(Tensors& t)
47+{
48+ const std::vector<float>* hosts[kInputCount] = {&kGwr, &kGwi, &kUr, &kUi, &kVr, &kVi};
49+ for (size_t i = 0; i < kInputCount; ++i) {
50+ auto ret = CreateAclTensor(*hosts[i], kShape, &t.inDev[i], aclDataType::ACL_FLOAT, &t.inT[i]);
51+ CHECK_RET(ret == ACL_SUCCESS, return ret);
52+ }
53+ const std::vector<float> packed(kPackedNumel, 0.0f);
54+ auto ret = CreateAclTensor(packed, kShape, &t.packedDev, aclDataType::ACL_FLOAT, &t.outT[0], kPackedNumel);
55+ CHECK_RET(ret == ACL_SUCCESS, return ret);
56+ for (size_t i = 1; i < kOutputCount; ++i) {
57+ ret = CreateAclTensorView(t.packedDev, kShape, aclDataType::ACL_FLOAT, &t.outT[i]);
58+ CHECK_RET(ret == ACL_SUCCESS, return ret);
59+ }
60+ return ACL_SUCCESS;
61+}
62+ 
63+int LaunchOp(const Tensors& t, aclrtStream stream, void** workspaceAddr)
64+{
65+ uint64_t workspaceSize = 0;
66+ aclOpExecutor* executor = nullptr;
67+ auto ret = aclnnComplexMulGradGetWorkspaceSize(t.inT[0], t.inT[1], t.inT[2], t.inT[3], t.inT[4], t.inT[5],
68+ t.outT[0], t.outT[1], t.outT[2], t.outT[3], &workspaceSize,
69+ &executor);
70+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnComplexMulGradGetWorkspaceSize failed. ERROR: %d\n", ret);
71+ return ret);
72+ ret = MallocWorkspace(workspaceSize, workspaceAddr);
73+ CHECK_RET(ret == ACL_SUCCESS, return ret);
74+ ret = aclnnComplexMulGrad(*workspaceAddr, workspaceSize, executor, stream);
75+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnComplexMulGrad failed. ERROR: %d\n", ret); return ret);
76+ ret = aclrtSynchronizeStream(stream);
77+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
78+ return ACL_SUCCESS;
79+}
80+ 
81+std::vector<float> Expected(const Tensors&)
82+{
83+ std::vector<float> expected(kPackedNumel, 0.0f);
84+ for (size_t i = 0; i < kNumel; ++i) {
85+ expected[i] = kGwr[i] * kVr[i] + kGwi[i] * kVi[i];
86+ expected[i + kNumel] = -kGwr[i] * kVi[i] + kGwi[i] * kVr[i];
87+ expected[i + kNumel * 2] = kGwr[i] * kUr[i] + kGwi[i] * kUi[i];
88+ expected[i + kNumel * 3] = -kGwr[i] * kUi[i] + kGwi[i] * kUr[i];
89+ }
90+ return expected;
91+}
92+ 
93+void Release(Tensors& t)
94+{
95+ DestroyTensors(t.inT, kInputCount);
96+ DestroyTensors(t.outT, kOutputCount);
97+ FreePackedBuffers(t.inDev, kInputCount, &t.packedDev);
98+}
99+ 
100+} // namespace
101+ 
102+int main()
103+{
104+ const AclnnExampleOps<Tensors> ops{SetupTensors, LaunchOp, &Tensors::packedDev, Expected, Release};
105+ return RunAclnnExample(ops, kPackedNumel);
106+}
Rtests/ascendc/ophost/test_complex_mul_tiling.cpp→tests/ut/op_host/complex_mul/test_complex_mul_tiling.cpp+5-59
Rtests/ascendc/ophost/test_fused_bias_softmax_tiling.cpp→tests/ut/op_host/fused_bias_softmax/test_fused_bias_softmax_tiling.cpp+15-2
Rtests/ascendc/ophost/test_fused_softmax_grad_tiling.cpp→tests/ut/op_host/fused_softmax_grad/test_fused_softmax_grad_tiling.cpp+2-1