oil-gas-ops-prospect算子开发指南从零新增一个AscendC算子的9步完整流程【免费下载链接】oil-gas-ops-prospect面向油气勘探oil gas exploration领域的昇腾自定义算子库。仓名即 oil-gas-ops油气算子 prospect勘探目标面向地震成像、全波形反演等勘探计算场景项目地址: https://gitcode.com/cann/oil-gas-ops-prospect本文带你从零开始在面向油气勘探领域的昇腾算子库oil-gas-ops-prospect中新增一个AscendC 算子。我们按「目录规划 → OpDef 定义 → tiling 切分 → kernel 实现 → manifest 注册 → Python 封装 → 单元测试 → 编译打包 → 910B 真机精度验收」共9 步拆解完整流程并给出每步的关键约定、命名规范和可参考的现成算子文件帮助新手少走弯路。一、开发前准备选对后端配好环境oil-gas-ops-prospect 是一个交付自定义算子的 CANN 仓库按交付物形态分为三条后端路线动手前先判断你的算子属于哪一类场景选择目录需要精细控制核间切分、UB 分配与流水追求极致性能AscendC本文主线ascendc/用 Python 快速表达、依赖 autograd、形态多变Triton-Ascendtriton-ascend/组合现有 torch 算子即可表达无需自定义 kernelComposite-Torchcomposite-torch/环境要求CANN 9.0.0含 Toolkit 对应 SoC 的 Ops 包 Triton-Ascend当前面向Ascend 910BAscend910B1。克隆仓库后安装 CANN 环境git clone https://gitcode.com/cann/oil-gas-ops-prospect.git完整安装步骤见 docs/zh/context/quick_install.md。二、第 1 步规划算子目录结构AscendC 算子统一放在ascendc/operators/op_name/下只放实现测试和样例一律放仓库根的tests/、examples/。以复数乘complex_mul为参考模板ascendc/operators/${op_name}/ ├── op_host/ │ ├── ${op_name}.cpp # OpDef、InferShape、TilingFunc │ ├── ${op_name}_tiling_data_def.h # TilingData 结构体定义 │ ├── ${op_name}_tiling_core.h # 平台无关的切分算法header-only │ └── ${op_name}_tiling_impl.h # 对接 gert::TilingContext 的实现 └── op_kernel/ ├── ${op_name}.cpp # kernel 入口 ├── ${op_name}_impl.h # kernel 实现 └── ${op_name}_tiling.h # kernel 侧 TilingData 视图 多算子共用的 host / kernel 头文件放 ascendc/operators/_common/合并进 msopgen 工程时会平铺到op_host/与op_kernel/。逐目录职责见 docs/zh/context/dir_structure.md。三、第 2 步编写 OpDef 与形状推理在op_host/${op_name}.cpp中用ops::OpDefOP_ADD声明算子注意不要用旧的REG_OP。三个必须实现的钩子OpDef声明输入/输出、dtype、format。逐元素 float32/ND 算子可直接复用 og_tiling_context.h 的DeclareRequiredFloatNdIo与AddAscend910bElementwiseConfigInferShape / InferDataType输出形状通常与输入一致可复用InferSameShapeOutputsTilingFunc挂到AICore().SetTiling(...)。另外两个关键约定opFile.value/opInterface.value必须等于op_kernel/下入口文件名无扩展名蛇形命名需要 aclnn 接口时设aclnnSupport.value support_aclnn。可参考的完整实现ascendc/operators/complex_mul/op_host/complex_mul.cpp。四、第 3 步编写可测的 tiling 切分算法⚠️ 这是新手最容易踩坑的一步。TilingFunc的入口签名依赖gert::TilingContextCANN 二进制包未提供 faker无法直接单测。因此本仓的约定是把与平台无关的切分算法抽到op_host/op_tiling_core.h单元测试直接驱动这个核心函数覆盖真实算法。TilingFunc本身只做四件事形状/dtype 校验 → 调*_tiling_core.h→SetBlockDim/SetTilingKey→SaveToBuffer。切分策略的三条经验法则参考 complex_mul_tiling_core.h按元素总数在 AIV 核间均分32B 对齐核内再按UB 预算分拍如前向 13 份 tile、反向 21 份GetCoreNumAiv()为 0 时必须失败或给出明确兜底并在 UT 中覆盖该分支。 硬件粒度提醒Ascend 910B 向量指令以64 个元素为粒度count不是 64 倍数时尾部可能被跳过。计算用 aligned 长度写出用原始count禁止越界写。五、第 4 步实现 kernel 入口op_kernel/${op_name}.cpp中定义__global__ __aicore__入口函数入口符号必须与opFile.value一致。典型流程拷贝 TilingData → 按TILING_KEY分发 →TPipe流水 算子类Init/Process。参考 AIV-only double buffer 的完整写法ascendc/operators/complex_mul/op_kernel/complex_mul.cpp。六、第 5 步注册算子并遵守命名约定在 ascendc/operators/manifest.tsv 追加一行tab 分隔目录名、OpType、kernel 入口文件名然后用cd ascendc bash build.sh list-ops确认注册成功。本仓有一套严格的命名约定务必照抄侧规则示例C OpTypePascalCaseComplexMulaclnnaclnn OpTypeaclnnComplexMulkernel 入口 /opFile.value蛇形与源文件名一致complex_mulPython 公开函数蛇形不加npu_前缀complex_mul反向 OPP 算子Grad后缀ComplexMulGrad环境变量OIL_GAS_OPS_前缀OIL_GAS_OPS_COMPLEX_MUL_IMPL 小技巧OpType 改名不影响 kernel 源码入口由opFile.value固定但 aclnn 符号会随 OpType 变化。七、第 6 步Python 封装与 aclnn 绑定Python 侧交付两个文件最终打包为oil_gas_ops_prospect.op_family模块文件作用ascendc/pybind/op_family.py公开 API直接成为模块ascendc/pybind/csrc/op_family/bindings.cpp直调 aclnn 两段式入口自动编成op_family_npu扩展新建csrc/op_family/*.cpp后会被pybind_extensions.py自动带上不用手写 setup.py。封装写法以 ascendc/pybind/complex_mul.py 为准规则如下函数式 API 走autograd.Function.apply直接调ext.forward()会断梯度图必须提供独立 PyTorch 参考实现对拍基准实现开关OIL_GAS_OPS_OP_IMPLtorch|ascendc|auto默认torchOIL_GAS_OPS_REQUIRE_ASCENDC_OP1时禁止静默回退扩展模块用lru_cache缓存 import。更多说明见 ascendc/pybind/README.md。八、第 7 步补充单元测试C 测试全部放在仓库根tests/算子目录不再放测试。三层覆盖tests/ut/op_host/${op_name}/test_${op_name}_tiling.cpp # tiling UTgoogletest tests/st/aclnnOpType/ # 多 case 两段式 aclnn 真值需 NPU examples/aclnn/${op_name}/test_aclnn_${op_name}.cpp # 内联小 tensor 样例需 NPUtests/ascendc/ophost/CMakeLists.txt会自动收集tests/ut/op_host/*/test_*.cpp无需改 CMakeST / 样例需登记到 tests/suites.json供build.sh -u编排。ophost UT 建议覆盖UB 预算约束、对齐与上限、小 tensor 收缩、退化输入负元素数、aivNum0、空张量、核间切分覆盖性不变量每个元素恰好被一个 AIV 核处理一次、区间无空洞无重叠。参考 tests/ut/op_host/complex_mul/test_complex_mul_tiling.cpp 与 tests/st/aclnnComplexMul/。九、第 8 步编译打包与安装仓库根目录build.sh是统一入口三档用法bash build.sh --pkg # 只出 OPP .run 包 → output/SOC/ bash build.sh # 编 .run 与 whl不安装 bash build.sh --install # 编译并安装到 CANN / Python并跑冒烟测试只构建单个新算子时加--ops${op_name}。产物统一落在output/安装后新 shell 需保留export LD_LIBRARY_PATH${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:$LD_LIBRARY_PATH完整参数见 docs/zh/context/build.md入口说明同 ascendc/README.md。十、第 9 步910B 真机精度与性能验收算子的精度、性能、显存必须在真实 Ascend 910B 上验证通过。在tests/ascendc/${op_name}/下按现有结构组织如 tests/ascendc/complex_mul/test_${op_name}_npu.pyrun_test.shREADME.md统一以等价的 PyTorch/torch_npu 表达式为基准。验收阈值参考已交付算子实测类型dtype阈值逐点乘类fp32相对误差≤ 1e-5融合 softmax 类fp32max_diff ≤ 1e-5融合 softmax 类bf16atol2e-2, rtol4e-3在 PR 中按 docs/zh/develop/precision_acceptance_template.md 填写精度验收报告通过率必须 100%严重/主要遗留问题未闭环不得合入。提交前自检清单bash build.sh -u # 全部测试套件需 910B bash scripts/check_license_header.sh # 许可证头检查Apache 2.0同时把算子语义、约束、签名同步写入 docs/zh/op_list.md、docs/zh/api_list.md 与ascendc/README.md。✅十一、常见问题 FAQQ1测试文件可以放在operators/op/tests/下吗不可以。operators/只保留op_host/op_kernelC 测试放仓库根tests/、样例放examples/aclnn/。Q2从 CANN 原仓移植的算子依赖内部框架头tiling_base.h、ops_err.h等怎么办提供本地 compat shim 替换OP_CHECK_IF、OPS_REPORT_*等宏只替换 include 不改动 tiling 算法确需额外头目录时在ascendc/scripts/patch_op_host.py的_EXTRA_INCLUDE_SUBDIRS中追加。详见 ascendc/README.md。Q3新增 SOC如 910C要做什么同时写入SUPPORTED_SOC_VERSIONS、在 OpDef 的AICore().AddConfig(...)增加配置并在该 SOC 上完成精度验收。十二、交付物速查表类别路径Host / Kernelascendc/operators/op/op_host/、op_kernel/含*_tiling_core.h注册表ascendc/operators/manifest.tsv 新行Python APIascendc/pybind/family.pyaclnn 绑定ascendc/pybind/csrc/family/bindings.cppophost UTtests/ut/op_host/op/test_op_tiling.cppaclnn 样例 / STexamples/aclnn/op/、tests/st/aclnnOpType/登记 tests/suites.jsonNPU 对比tests/ascendc/op/文档docs/zh/op_list.md、docs/zh/api_list.md、ascendc/README.md精度报告PR 附 验收模板掌握以上 9 步你就拥有了在 oil-gas-ops-prospect 中新增 AscendC 算子的完整路线图。更多细节请阅读官方开发指南 docs/zh/develop/operator_development_guide.md并对照complex_mul/fused_bias_softmax等现成算子动手实践吧【免费下载链接】oil-gas-ops-prospect面向油气勘探oil gas exploration领域的昇腾自定义算子库。仓名即 oil-gas-ops油气算子 prospect勘探目标面向地震成像、全波形反演等勘探计算场景项目地址: https://gitcode.com/cann/oil-gas-ops-prospect创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
oil-gas-ops-prospect算子开发指南：从零新增一个AscendC算子的9步完整流程
oil-gas-ops-prospect算子开发指南从零新增一个AscendC算子的9步完整流程【免费下载链接】oil-gas-ops-prospect面向油气勘探oil gas exploration领域的昇腾自定义算子库。仓名即 oil-gas-ops油气算子 prospect勘探目标面向地震成像、全波形反演等勘探计算场景项目地址: https://gitcode.com/cann/oil-gas-ops-prospect本文带你从零开始在面向油气勘探领域的昇腾算子库oil-gas-ops-prospect中新增一个AscendC 算子。我们按「目录规划 → OpDef 定义 → tiling 切分 → kernel 实现 → manifest 注册 → Python 封装 → 单元测试 → 编译打包 → 910B 真机精度验收」共9 步拆解完整流程并给出每步的关键约定、命名规范和可参考的现成算子文件帮助新手少走弯路。一、开发前准备选对后端配好环境oil-gas-ops-prospect 是一个交付自定义算子的 CANN 仓库按交付物形态分为三条后端路线动手前先判断你的算子属于哪一类场景选择目录需要精细控制核间切分、UB 分配与流水追求极致性能AscendC本文主线ascendc/用 Python 快速表达、依赖 autograd、形态多变Triton-Ascendtriton-ascend/组合现有 torch 算子即可表达无需自定义 kernelComposite-Torchcomposite-torch/环境要求CANN 9.0.0含 Toolkit 对应 SoC 的 Ops 包 Triton-Ascend当前面向Ascend 910BAscend910B1。克隆仓库后安装 CANN 环境git clone https://gitcode.com/cann/oil-gas-ops-prospect.git完整安装步骤见 docs/zh/context/quick_install.md。二、第 1 步规划算子目录结构AscendC 算子统一放在ascendc/operators/op_name/下只放实现测试和样例一律放仓库根的tests/、examples/。以复数乘complex_mul为参考模板ascendc/operators/${op_name}/ ├── op_host/ │ ├── ${op_name}.cpp # OpDef、InferShape、TilingFunc │ ├── ${op_name}_tiling_data_def.h # TilingData 结构体定义 │ ├── ${op_name}_tiling_core.h # 平台无关的切分算法header-only │ └── ${op_name}_tiling_impl.h # 对接 gert::TilingContext 的实现 └── op_kernel/ ├── ${op_name}.cpp # kernel 入口 ├── ${op_name}_impl.h # kernel 实现 └── ${op_name}_tiling.h # kernel 侧 TilingData 视图 多算子共用的 host / kernel 头文件放 ascendc/operators/_common/合并进 msopgen 工程时会平铺到op_host/与op_kernel/。逐目录职责见 docs/zh/context/dir_structure.md。三、第 2 步编写 OpDef 与形状推理在op_host/${op_name}.cpp中用ops::OpDefOP_ADD声明算子注意不要用旧的REG_OP。三个必须实现的钩子OpDef声明输入/输出、dtype、format。逐元素 float32/ND 算子可直接复用 og_tiling_context.h 的DeclareRequiredFloatNdIo与AddAscend910bElementwiseConfigInferShape / InferDataType输出形状通常与输入一致可复用InferSameShapeOutputsTilingFunc挂到AICore().SetTiling(...)。另外两个关键约定opFile.value/opInterface.value必须等于op_kernel/下入口文件名无扩展名蛇形命名需要 aclnn 接口时设aclnnSupport.value support_aclnn。可参考的完整实现ascendc/operators/complex_mul/op_host/complex_mul.cpp。四、第 3 步编写可测的 tiling 切分算法⚠️ 这是新手最容易踩坑的一步。TilingFunc的入口签名依赖gert::TilingContextCANN 二进制包未提供 faker无法直接单测。因此本仓的约定是把与平台无关的切分算法抽到op_host/op_tiling_core.h单元测试直接驱动这个核心函数覆盖真实算法。TilingFunc本身只做四件事形状/dtype 校验 → 调*_tiling_core.h→SetBlockDim/SetTilingKey→SaveToBuffer。切分策略的三条经验法则参考 complex_mul_tiling_core.h按元素总数在 AIV 核间均分32B 对齐核内再按UB 预算分拍如前向 13 份 tile、反向 21 份GetCoreNumAiv()为 0 时必须失败或给出明确兜底并在 UT 中覆盖该分支。 硬件粒度提醒Ascend 910B 向量指令以64 个元素为粒度count不是 64 倍数时尾部可能被跳过。计算用 aligned 长度写出用原始count禁止越界写。五、第 4 步实现 kernel 入口op_kernel/${op_name}.cpp中定义__global__ __aicore__入口函数入口符号必须与opFile.value一致。典型流程拷贝 TilingData → 按TILING_KEY分发 →TPipe流水 算子类Init/Process。参考 AIV-only double buffer 的完整写法ascendc/operators/complex_mul/op_kernel/complex_mul.cpp。六、第 5 步注册算子并遵守命名约定在 ascendc/operators/manifest.tsv 追加一行tab 分隔目录名、OpType、kernel 入口文件名然后用cd ascendc bash build.sh list-ops确认注册成功。本仓有一套严格的命名约定务必照抄侧规则示例C OpTypePascalCaseComplexMulaclnnaclnn OpTypeaclnnComplexMulkernel 入口 /opFile.value蛇形与源文件名一致complex_mulPython 公开函数蛇形不加npu_前缀complex_mul反向 OPP 算子Grad后缀ComplexMulGrad环境变量OIL_GAS_OPS_前缀OIL_GAS_OPS_COMPLEX_MUL_IMPL 小技巧OpType 改名不影响 kernel 源码入口由opFile.value固定但 aclnn 符号会随 OpType 变化。七、第 6 步Python 封装与 aclnn 绑定Python 侧交付两个文件最终打包为oil_gas_ops_prospect.op_family模块文件作用ascendc/pybind/op_family.py公开 API直接成为模块ascendc/pybind/csrc/op_family/bindings.cpp直调 aclnn 两段式入口自动编成op_family_npu扩展新建csrc/op_family/*.cpp后会被pybind_extensions.py自动带上不用手写 setup.py。封装写法以 ascendc/pybind/complex_mul.py 为准规则如下函数式 API 走autograd.Function.apply直接调ext.forward()会断梯度图必须提供独立 PyTorch 参考实现对拍基准实现开关OIL_GAS_OPS_OP_IMPLtorch|ascendc|auto默认torchOIL_GAS_OPS_REQUIRE_ASCENDC_OP1时禁止静默回退扩展模块用lru_cache缓存 import。更多说明见 ascendc/pybind/README.md。八、第 7 步补充单元测试C 测试全部放在仓库根tests/算子目录不再放测试。三层覆盖tests/ut/op_host/${op_name}/test_${op_name}_tiling.cpp # tiling UTgoogletest tests/st/aclnnOpType/ # 多 case 两段式 aclnn 真值需 NPU examples/aclnn/${op_name}/test_aclnn_${op_name}.cpp # 内联小 tensor 样例需 NPUtests/ascendc/ophost/CMakeLists.txt会自动收集tests/ut/op_host/*/test_*.cpp无需改 CMakeST / 样例需登记到 tests/suites.json供build.sh -u编排。ophost UT 建议覆盖UB 预算约束、对齐与上限、小 tensor 收缩、退化输入负元素数、aivNum0、空张量、核间切分覆盖性不变量每个元素恰好被一个 AIV 核处理一次、区间无空洞无重叠。参考 tests/ut/op_host/complex_mul/test_complex_mul_tiling.cpp 与 tests/st/aclnnComplexMul/。九、第 8 步编译打包与安装仓库根目录build.sh是统一入口三档用法bash build.sh --pkg # 只出 OPP .run 包 → output/SOC/ bash build.sh # 编 .run 与 whl不安装 bash build.sh --install # 编译并安装到 CANN / Python并跑冒烟测试只构建单个新算子时加--ops${op_name}。产物统一落在output/安装后新 shell 需保留export LD_LIBRARY_PATH${ASCEND_HOME_PATH}/opp/vendors/customize/op_api/lib:${ASCEND_HOME_PATH}/lib64:$LD_LIBRARY_PATH完整参数见 docs/zh/context/build.md入口说明同 ascendc/README.md。十、第 9 步910B 真机精度与性能验收算子的精度、性能、显存必须在真实 Ascend 910B 上验证通过。在tests/ascendc/${op_name}/下按现有结构组织如 tests/ascendc/complex_mul/test_${op_name}_npu.pyrun_test.shREADME.md统一以等价的 PyTorch/torch_npu 表达式为基准。验收阈值参考已交付算子实测类型dtype阈值逐点乘类fp32相对误差≤ 1e-5融合 softmax 类fp32max_diff ≤ 1e-5融合 softmax 类bf16atol2e-2, rtol4e-3在 PR 中按 docs/zh/develop/precision_acceptance_template.md 填写精度验收报告通过率必须 100%严重/主要遗留问题未闭环不得合入。提交前自检清单bash build.sh -u # 全部测试套件需 910B bash scripts/check_license_header.sh # 许可证头检查Apache 2.0同时把算子语义、约束、签名同步写入 docs/zh/op_list.md、docs/zh/api_list.md 与ascendc/README.md。✅十一、常见问题 FAQQ1测试文件可以放在operators/op/tests/下吗不可以。operators/只保留op_host/op_kernelC 测试放仓库根tests/、样例放examples/aclnn/。Q2从 CANN 原仓移植的算子依赖内部框架头tiling_base.h、ops_err.h等怎么办提供本地 compat shim 替换OP_CHECK_IF、OPS_REPORT_*等宏只替换 include 不改动 tiling 算法确需额外头目录时在ascendc/scripts/patch_op_host.py的_EXTRA_INCLUDE_SUBDIRS中追加。详见 ascendc/README.md。Q3新增 SOC如 910C要做什么同时写入SUPPORTED_SOC_VERSIONS、在 OpDef 的AICore().AddConfig(...)增加配置并在该 SOC 上完成精度验收。十二、交付物速查表类别路径Host / Kernelascendc/operators/op/op_host/、op_kernel/含*_tiling_core.h注册表ascendc/operators/manifest.tsv 新行Python APIascendc/pybind/family.pyaclnn 绑定ascendc/pybind/csrc/family/bindings.cppophost UTtests/ut/op_host/op/test_op_tiling.cppaclnn 样例 / STexamples/aclnn/op/、tests/st/aclnnOpType/登记 tests/suites.jsonNPU 对比tests/ascendc/op/文档docs/zh/op_list.md、docs/zh/api_list.md、ascendc/README.md精度报告PR 附 验收模板掌握以上 9 步你就拥有了在 oil-gas-ops-prospect 中新增 AscendC 算子的完整路线图。更多细节请阅读官方开发指南 docs/zh/develop/operator_development_guide.md并对照complex_mul/fused_bias_softmax等现成算子动手实践吧【免费下载链接】oil-gas-ops-prospect面向油气勘探oil gas exploration领域的昇腾自定义算子库。仓名即 oil-gas-ops油气算子 prospect勘探目标面向地震成像、全波形反演等勘探计算场景项目地址: https://gitcode.com/cann/oil-gas-ops-prospect创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
