CANN 鸿蒙端侧 Sobel 边缘检测自定义算子开发指南基于 Ascend C 的 CV 前处理 NPU 迁移实践【免费下载链接】cann-recipes-harmony-infer本项目为鸿蒙开发者提供基于CANN平台的业务实践案例方便开发者参考实现端云能力迁移及端侧推理部署。项目地址: https://gitcode.com/cann/cann-recipes-harmony-infer导读本文面向鸿蒙开发者与昇腾算子开发工程师讲解如何在 CANN 平台上使用 Ascend C 编程语言实现 Sobel 边缘检测自定义算子SobelCustom将传统上依赖 CPU 的图像 CV 前处理搬移到 NPU 上计算。文中完整覆盖算子数学原理、算子工程结构、Host 侧 Tiling 与 Kernel 侧实现细节、ONNX 框架插件适配、编译安装、golden 数据验证以及基于 HiAI Foundation / NNCore 的鸿蒙端到端模型部署链路读者可据此复现工程创建 → 算子实现 → 编译部署 → 算子调用的完整闭环。鸿蒙端 SobelCustom 应用主界面原始图像来自 Soble 应用截图NPU 推理完成后的边缘检测结果显示 NPU 运行时间来自 Soble 应用截图一、应用场景为什么要把 CV 前处理搬到 NPU图像 CV 前处理的计算逻辑灵活多变常规 NN 模型较难处理也难以直接使用 NPU 上的常规 NN 算子进行表达。为了在 NPU 上解决 CV 前处理灵活计算的问题CANN 与某伙伴合作使用 Ascend C 实现了 CV 领域的 Sobel 计算将原本在 CPU 侧串行执行的前处理步骤灰度化、卷积、梯度融合整体搬移到 NPU 上并行执行。该方案已在某客户的边缘检测业务中得到验证和应用。这一迁移的价值在于端侧推理流水线中前处理往往是串行瓶颈把 Sobel 这类计算规则明确、数据局部性强的算法固化为自定义算子后前处理与模型推理可以在同一 NPU 流水上完成减少 CPU 与 NPU 间的数据搬运与同步开销。二、SobelCustom 算子样例说明与开发流程本样例通过 Ascend C 编程语言实现 Sobel 算子。整个开发过程遵循昇腾自定义算子开发的标准流程工程创建基于算子工程框架CMake 工程 目录骨架搭建工程算子实现完成算子核函数Kernel开发和 Tiling 实现Host 侧编译部署通过工程编译脚本完成算子的编译与安装生成自定义算子 Run 包算子调用实现单算子调用或在第三方框架ONNX → ATC → omc中完成算子调用最终在鸿蒙端侧通过模型推理接口执行。整个过程都依赖算子工程工程代码框架承载核函数与 Tiling 实现工程编译脚本负责编译部署产物对接单算子调用或第三方框架中的算子调用。三、算子描述与数学原理Sobel 算子是图像边缘检测的经典一阶微分算子通过两个 3×3 卷积核分别计算水平方向Sobel-x与垂直方向Sobel-y的梯度再对二者取绝对值求和得到边缘强度图。算子对应的数学表达式为dx[i, j] A[idi, jdj] * Cx[di, dj], di [-1, 0, 1], dj [-1, 0, 1]其中A为灰度化后的输入图像Cx为水平 Sobel 卷积核同理可构造dy与垂直核Cy。最终输出为abs(dx) abs(dy)再经取整与饱和裁剪0~255得到 uint8 边缘图。结合仓库中 gen_data.py 的 golden 数据生成逻辑可以确认算子的完整计算链RGB2Gray0.299 * r 0.587 * g 0.114 * bSobel-x / Sobel-y分别以 3×3 卷积核计算水平、垂直梯度并取绝对值np.abs(sobel_x)梯度融合y2 |dx| |dy|量化输出np.clip(y2, 0, 255)后np.ceil并转 uint8。四、算子规格描述本样例支持如下输入规格输入张量1 × 763 × 1024 × 3batch1H763W1024C3数据类型uint8ND 布局NHWC输出张量1 × 1 × 761 × 1022数据类型uint8。输出尺寸比输入在 H、W 两个维度各减 23×3 卷积核不补零的有效卷积结果通道数由 3 变为 1RGB 融合为单通道灰度梯度图。上述规格在 create_onnx.py 中定义并由 sobel_custom.cppop_host 的InferShape逻辑推导得出int output_height x1_shape-GetDim(1) - 2; // 减去卷积核的高度减2 int output_width x1_shape-GetDim(2) - 2; // 减去卷积核的宽度减2 // y_shape: [batch, channels(即输入C-2), output_height, output_width] y_shape-SetDim(1, x1_shape-GetDim(3) - 2); // channels y_shape-SetDim(2, output_height); y_shape-SetDim(3, output_width);五、支持的产品型号本样例支持如下产品型号对应 op_host/sobel_custom.cpp 中的AddConfig(kirinx90, ...)配置项Kirin X90 处理器系列产品Kirin 9030 处理器系列产品工程根目录的 CMakePresets.json 中通过ASCEND_COMPUTE_UNIT: kirinx90指定目标计算单元vendor_name为customize。六、算子工程结构SobelCustom 算子工程位于 ops/ascendc/src/sobel_custom采用昇腾算子工程标准三段式结构sobel_custom/ ├── CMakeLists.txt # 工程顶层构建脚本npu_op_package ├── CMakePresets.json # CMake 预置配置ASCEND_COMPUTE_UNIT 等 ├── build_and_install.sh # 编译安装脚本生成并安装 custom_opp 包 ├── framework/ │ └── onnx_plugin/ │ └── sobel_custom_plugin.cc # ONNX 框架算子插件注册 ├── op_host/ │ ├── sobel_custom.cpp # OpDef 定义、InferShape、TilingFunc │ └── sobel_custom_tiling.h # TilingData 结构定义与注册 ├── op_kernel/ │ ├── sobel_custom.cpp # 核函数CopyIn/Compute/CopyOut │ └── sobel_custom_base.h # Min / CeilDiv 工具函数 └── test/ ├── create_onnx.py # 生成单算子 ONNX 模型 └── gen_data.py # 生成输入 x.bin 与 golden y.bin顶层 CMakeLists.txt 通过npu_op_package打包算子并分别add_subdirectory(framework)、op_host、op_kernel。七、算子实现Host 侧算子定义与 Tiling7.1 算子定义OpDef在 op_host/sobel_custom.cpp 中SobelCustom类继承OpDef完成算子注册输入x必选参数数据类型ge::DT_UINT8格式FORMAT_ND输出y必选参数数据类型ge::DT_UINT8格式FORMAT_ND绑定InferShape、InferDataType回调输出数据类型与输入一致AICore 配置开启动态编译、动态 Format、动态 Rank、动态 Shape 支持通过AddConfig(kirinx90, aicore_config)注册 Kirin X90 处理器配置。7.2 Tiling 实现Tiling 策略在 sobel_custom.cpp 的 TilingFunc 中计算核心参数如下常量值含义USE_CORE_NUM1使用 1 个 AI CoreBUFFER_NUM2双缓冲队列深度tileNum8循环分块数量h9Tile 高含 3×3 卷积所需的 halo 重叠行w256Tile 宽c3通道数Tiling 计算逻辑由输入GetShapeSize()得到totalLengthblockLength totalLength / USE_CORE_NUMtileLength h * w * cgrayLength h * w并读取输入 Shape 的 H、W 维度写入 Tiling 数据最后context-SetBlockDim(1)设定单核执行Workspace 大小为 0。7.3 TilingData 结构sobel_custom_tiling.h 通过BEGIN_TILING_DATA_DEF / END_TILING_DATA_DEF宏定义并在REGISTER_TILING_DATA_CLASS(SobelCustom, SobelCustomTilingData)注册TILING_DATA_FIELD_DEF(uint32_t, size); TILING_DATA_FIELD_DEF(uint32_t, totalLength); TILING_DATA_FIELD_DEF(uint32_t, blockLength); TILING_DATA_FIELD_DEF(uint32_t, tileLength); TILING_DATA_FIELD_DEF(uint32_t, grayLength); TILING_DATA_FIELD_DEF(uint32_t, tileNum); TILING_DATA_FIELD_DEF(uint32_t, H); TILING_DATA_FIELD_DEF(uint32_t, W);这些字段在 Kernel 侧通过GET_TILING_DATA取出用于驱动核函数的搬运与计算。八、算子实现Kernel 侧CopyIn / Compute / CopyOut核函数主体位于 op_kernel/sobel_custom.cpp模板类KernelSobelCustomuint8_t采用 Ascend C 经典的搬入-计算-搬出三段流水结构配合双缓冲队列TQueQuePosition::VECIN, BUFFER_NUM隐藏搬运与计算延迟。8.1 数据分块与搬入CopyInProcess()中按cntH CeilDiv(H, h)、cntW CeilDiv(W, w)将输入切分为h × w的 Tile逐块执行CopyIn → Compute → CopyOut。CopyIn针对不同位置首列、中间块、末列、最后一行分别构造DataCopyExtParams通过srcStride/dstStride处理行间跳步并使用DataCopyPad统一搬运带 halo 重叠3×3 卷积需要相邻 Tile 共享边界行/列的数据块。CeilDiv与Min工具函数定义在 sobel_custom_base.h。8.2 计算流程ComputeCompute()是算法核心代码 中的计算链如下NHWC → NCHW 转置AscendC::TransposeTRANSPOSE_NHWC2NCHW将输入转换为通道分离布局便于逐通道处理类型提升AscendC::Cast将 uint8 转为 halfFP16提升中间计算精度RGB2Gray利用Muls按0.299 / 0.587 / 0.114加权三个通道后Add合并得到灰度图Sobel-x 计算CreateVecIndex生成索引序列[2, 3, ..., w-1]Muls缩放后Cast为 uint32由于data[i][j1]、data[i][j2]等偏移地址不满足 32 字节对齐代码使用AscendC::Gather按索引取数再以Muls、Add完成-1·P0 - 2·P1 - 1·P2 2·P4 1·P5 2·P6 1·P7的卷积累加Sobel-y 计算同理生成两套索引[1..w-2]与[2..w-1]通过Gather处理非对齐访问完成垂直方向梯度梯度融合Abs(dx)、Abs(dy)取绝对值后Add得到|dx| |dy|量化输出Cast(..., CAST_CEIL)将 half 结果向上取整转回 uint8。代码注释明确说明了使用Gather的原因由于data[i][j 2]的地址不是 32Byte 对齐的所以需要用 gather这正是向量指令对地址对齐要求的典型处理手法。8.3 结果搬出CopyOutCopyOut按h-2 × w-2的有效输出块剔除 halo 行/列组织DataCopyExtParams对不同边界块左上、最后一列、最后一行分别设置blockCount、blockLen与srcStride/dstStride最终DataCopyPad写回 Global Memory。8.4 核函数入口extern C __global__ __aicore__ void sobel_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tiling_data, tiling); KernelSobelCustomuint8_t op; op.Init(x, y, tiling_data); op.Process(); }九、ONNX 框架插件适配为使 ATC 工具能够将 ONNX 模型中的SobelCustom节点映射到本自定义算子需要在 framework/onnx_plugin/sobel_custom_plugin.cc 中完成 GE 侧注册REGISTER_CUSTOM_OP(SobelCustom) // 算子注册名 .FrameworkType(ONNX) // 原始框架类型 .OriginOpType(ai.onnx::11::SobelCustom) // 原始框架中的算子类型 .ParseParamsByOperatorFn(ParseParamSobelCustom);ParseParamSobelCustom当前直接返回SUCCESS该算子无额外属性需要解析。对应的单算子 ONNX 模型由 create_onnx.py 生成使用 ONNX Helper 构造输入[1, 763, 1024, 3]UINT8、输出[1, 1, 761, 1022]UINT8的SobelCustom节点opset 版本 11输出SobelCustom.onnx供 ATC 转换。十、编译安装与调试验证10.1 环境准备编译前请参考环境准备文档完成环境搭建确保开发套件 Ascend-cann-toolkit 安装完成并确认算子目录 CMakePresets.json 中ASCEND_CANN_PACKAGE_PATH字段设置为正确的 toolkit 安装路径一般为${install_path}/cann。10.2 编译安装在算子工程目录ops/ascendc/src/sobel_custom下执行# 为脚本添加执行权限 chmod x build_and_install.sh # 执行编译安装 ./build_and_install.shbuild_and_install.sh 的主要流程若未设置ASCEND_HOME_PATH默认取/usr/local/Ascend/cann并source其setenv.bash导出ASCEND_TENSOR_COMPILER_INCLUDE与LD_PRELOADlibmmpa.so按实际路径修正CMakePresets.json中的 CANN 路径以cmake --presetdefault配置随后依次构建binary与package目标在build_out下执行生成的custom_opp_${OS_ID}_${arch}.run --quiet完成算子包安装。更多编译细节可参考算子工程编译安装指南。10.3 单算子验证golden 数据test/gen_data.py 生成输入与参考输出固定随机种子np.random.seed(42)生成(1, 763, 1024, 3)的 uint8 随机图写入x.bin用 NumPy 按0.299r 0.587g 0.114b计算灰度图用 OpenCVcv2.Sobel(..., ksize3)分别计算 x、y 方向梯度取绝对值后求和、裁剪到[0, 255]、np.ceil转 uint8写入y.bin作为 golden 数据。该脚本同时扮演可执行规格说明的角色与 Kernel 侧 Compute() 的算子语义一一对应可用于校验 NPU 输出精度。10.4 调试与端到端转换执行test目录下的create_onnx.py与gen_data.py生成单算子 ONNX 模型、输入与 golden 数据后可通过 ATC 工具指南 将模型转换为鸿蒙端可加载的 omc 模型再调用鸿蒙维测接口完成单算子的性能与精度验证。十一、鸿蒙端端到端部署Soble 应用11.1 应用工程模型端到端部署参考 Soble 鸿蒙应用说明。应用工程位于 harmony_infer/harmony_os_next/Soble核心代码entry/src/main/ ├── cpp/ │ ├── SobelCustom.cpp // native 层接口实现 │ ├── HIAIModelManager.cpp // 模型管理类实现 │ ├── HIAIModelManager.h // 模型管理类定义 │ ├── CMakeLists.txt // native 层编译配置 │ └── types/libentry/ // native 接口注册 ├── ets/ │ ├── entryability/EntryAbility.ets │ └── pages/Index.ets // 主界面CPU/NPU 推理切换 └── resources/ ├── base/media/ // cup.jpg、guitar.jpg 等测试图片 └── rawfile/SobelCustom.omc // ATC 转换后的模型文件11.2 部署步骤使用 DevEco 构建应用前将 ATC 转换后的SobelCustom.omc模型文件放置到应用entry/src/main/resources/rawfile目录下再进行应用构建和安装在手机主屏幕点击NPU 推理应用自动加载模型并执行推理展示处理后的图片与模型处理时间运行截图见文首NPU 处理耗时以界面实测显示为准点击CPU 推理可将图片交由 CPU 处理便于对比 CPU 与 NPU 的处理时间点击Click for next image切换下一张测试图片退出应用时自动卸载模型。11.3 底层 API 调用链应用依赖 CANN 动态库libhiai_foundation.so与 AI 领域公共动态库libneural_network_core.sonative 侧通过 HiAI Foundation 与 NNCore 定义的一组 API 完成模型加载 → 编译 → 执行模型兼容性检查与加载HMS_HiAICompatibility_CheckFromBuffer、OH_NNCompilation_ConstructWithOfflineModelBuffer编译与设备设置OH_NNCompilation_SetDevice、HMS_HiAIOptions_SetBandMode、HMS_HiAIOptions_SetModelDeviceOrder、OH_NNCompilation_BuildExecutor 构造与张量管理OH_NNExecutor_Construct、OH_NNExecutor_CreateInputTensorDesc/CreateOutputTensorDesc、OH_NNTensor_Create、OH_NNTensor_GetDataBuffer/GetSize同步推理OH_NNExecutor_RunSync(executor, inputTensor[], inputCount, outputTensor[], outputCount)资源释放OH_NNTensor_Destroy、OH_NNTensorDesc_Destroy、OH_NNCompilation_Destroy、OH_NNExecutor_Destroy。OH_NNExecutor_RunSync是推理的关键入口将待处理的图片数据填充到输入 Tensor 后同步执行输出 Tensor 即 Sobel 边缘检测结果。该应用的模型管理实现可参考 HIAIModelManager.cpp界面逻辑参考 Index.ets。11.4 约束与限制本示例仅支持标准系统上运行支持华为手机、平板和 2in1 设备HarmonyOS 系统HarmonyOS 5.0.1 Release 及以上DevEco Studio 版本DevEco Studio 6.0.0 Release 及以上HarmonyOS SDK 版本HarmonyOS 6.0.0 Release SDK 及以上。十二、总结本样例完整展示了从 Ascend C 算子实现到鸿蒙端侧推理的整条链路算法侧Sobel 算子的RGB2Gray → Sobel-x/y → |dx||dy|全流程向量化通过Transpose、Cast、Gather、Muls、Add、Abs等 Ascend C 指令实现并正确处理了非 32 字节对齐访问工程侧标准三段式算子工程framework / op_host / op_kernel配合build_and_install.sh一键编译安装ONNX 插件注册打通 ATC 转换验证侧create_onnx.pygen_data.py提供可复现的单算子模型与 golden 数据部署侧ATC 产物接入鸿蒙 Soble 应用通过 HiAI Foundation / NNCore API 在 NPU 上完成推理实现 CV 前处理从 CPU 到 NPU 的迁移。对于图像类端侧推理场景SobelCustom 提供了一种可复用的范式将灵活多变的 CV 前处理以自定义算子形式固化到 NPU 执行流水从而减少 CPU-NPU 数据往返、提升端侧推理整体效率。相关文档与源码索引本文关联文档custom-npu_sobel.md算子工程ops/ascendc/src/sobel_custom编译安装指南ops/ascendc/src/README.md环境准备docs/quick_install.mdATC 工具docs/atc_tools_guide.md鸿蒙端部署Soble/readme_cn.md算子开发通用指南ascendc_develop_guide.md【免费下载链接】cann-recipes-harmony-infer本项目为鸿蒙开发者提供基于CANN平台的业务实践案例方便开发者参考实现端云能力迁移及端侧推理部署。项目地址: https://gitcode.com/cann/cann-recipes-harmony-infer创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN 鸿蒙端侧 Sobel 边缘检测自定义算子开发指南：基于 Ascend C 的 CV 前处理 NPU 迁移实践
CANN 鸿蒙端侧 Sobel 边缘检测自定义算子开发指南基于 Ascend C 的 CV 前处理 NPU 迁移实践【免费下载链接】cann-recipes-harmony-infer本项目为鸿蒙开发者提供基于CANN平台的业务实践案例方便开发者参考实现端云能力迁移及端侧推理部署。项目地址: https://gitcode.com/cann/cann-recipes-harmony-infer导读本文面向鸿蒙开发者与昇腾算子开发工程师讲解如何在 CANN 平台上使用 Ascend C 编程语言实现 Sobel 边缘检测自定义算子SobelCustom将传统上依赖 CPU 的图像 CV 前处理搬移到 NPU 上计算。文中完整覆盖算子数学原理、算子工程结构、Host 侧 Tiling 与 Kernel 侧实现细节、ONNX 框架插件适配、编译安装、golden 数据验证以及基于 HiAI Foundation / NNCore 的鸿蒙端到端模型部署链路读者可据此复现工程创建 → 算子实现 → 编译部署 → 算子调用的完整闭环。鸿蒙端 SobelCustom 应用主界面原始图像来自 Soble 应用截图NPU 推理完成后的边缘检测结果显示 NPU 运行时间来自 Soble 应用截图一、应用场景为什么要把 CV 前处理搬到 NPU图像 CV 前处理的计算逻辑灵活多变常规 NN 模型较难处理也难以直接使用 NPU 上的常规 NN 算子进行表达。为了在 NPU 上解决 CV 前处理灵活计算的问题CANN 与某伙伴合作使用 Ascend C 实现了 CV 领域的 Sobel 计算将原本在 CPU 侧串行执行的前处理步骤灰度化、卷积、梯度融合整体搬移到 NPU 上并行执行。该方案已在某客户的边缘检测业务中得到验证和应用。这一迁移的价值在于端侧推理流水线中前处理往往是串行瓶颈把 Sobel 这类计算规则明确、数据局部性强的算法固化为自定义算子后前处理与模型推理可以在同一 NPU 流水上完成减少 CPU 与 NPU 间的数据搬运与同步开销。二、SobelCustom 算子样例说明与开发流程本样例通过 Ascend C 编程语言实现 Sobel 算子。整个开发过程遵循昇腾自定义算子开发的标准流程工程创建基于算子工程框架CMake 工程 目录骨架搭建工程算子实现完成算子核函数Kernel开发和 Tiling 实现Host 侧编译部署通过工程编译脚本完成算子的编译与安装生成自定义算子 Run 包算子调用实现单算子调用或在第三方框架ONNX → ATC → omc中完成算子调用最终在鸿蒙端侧通过模型推理接口执行。整个过程都依赖算子工程工程代码框架承载核函数与 Tiling 实现工程编译脚本负责编译部署产物对接单算子调用或第三方框架中的算子调用。三、算子描述与数学原理Sobel 算子是图像边缘检测的经典一阶微分算子通过两个 3×3 卷积核分别计算水平方向Sobel-x与垂直方向Sobel-y的梯度再对二者取绝对值求和得到边缘强度图。算子对应的数学表达式为dx[i, j] A[idi, jdj] * Cx[di, dj], di [-1, 0, 1], dj [-1, 0, 1]其中A为灰度化后的输入图像Cx为水平 Sobel 卷积核同理可构造dy与垂直核Cy。最终输出为abs(dx) abs(dy)再经取整与饱和裁剪0~255得到 uint8 边缘图。结合仓库中 gen_data.py 的 golden 数据生成逻辑可以确认算子的完整计算链RGB2Gray0.299 * r 0.587 * g 0.114 * bSobel-x / Sobel-y分别以 3×3 卷积核计算水平、垂直梯度并取绝对值np.abs(sobel_x)梯度融合y2 |dx| |dy|量化输出np.clip(y2, 0, 255)后np.ceil并转 uint8。四、算子规格描述本样例支持如下输入规格输入张量1 × 763 × 1024 × 3batch1H763W1024C3数据类型uint8ND 布局NHWC输出张量1 × 1 × 761 × 1022数据类型uint8。输出尺寸比输入在 H、W 两个维度各减 23×3 卷积核不补零的有效卷积结果通道数由 3 变为 1RGB 融合为单通道灰度梯度图。上述规格在 create_onnx.py 中定义并由 sobel_custom.cppop_host 的InferShape逻辑推导得出int output_height x1_shape-GetDim(1) - 2; // 减去卷积核的高度减2 int output_width x1_shape-GetDim(2) - 2; // 减去卷积核的宽度减2 // y_shape: [batch, channels(即输入C-2), output_height, output_width] y_shape-SetDim(1, x1_shape-GetDim(3) - 2); // channels y_shape-SetDim(2, output_height); y_shape-SetDim(3, output_width);五、支持的产品型号本样例支持如下产品型号对应 op_host/sobel_custom.cpp 中的AddConfig(kirinx90, ...)配置项Kirin X90 处理器系列产品Kirin 9030 处理器系列产品工程根目录的 CMakePresets.json 中通过ASCEND_COMPUTE_UNIT: kirinx90指定目标计算单元vendor_name为customize。六、算子工程结构SobelCustom 算子工程位于 ops/ascendc/src/sobel_custom采用昇腾算子工程标准三段式结构sobel_custom/ ├── CMakeLists.txt # 工程顶层构建脚本npu_op_package ├── CMakePresets.json # CMake 预置配置ASCEND_COMPUTE_UNIT 等 ├── build_and_install.sh # 编译安装脚本生成并安装 custom_opp 包 ├── framework/ │ └── onnx_plugin/ │ └── sobel_custom_plugin.cc # ONNX 框架算子插件注册 ├── op_host/ │ ├── sobel_custom.cpp # OpDef 定义、InferShape、TilingFunc │ └── sobel_custom_tiling.h # TilingData 结构定义与注册 ├── op_kernel/ │ ├── sobel_custom.cpp # 核函数CopyIn/Compute/CopyOut │ └── sobel_custom_base.h # Min / CeilDiv 工具函数 └── test/ ├── create_onnx.py # 生成单算子 ONNX 模型 └── gen_data.py # 生成输入 x.bin 与 golden y.bin顶层 CMakeLists.txt 通过npu_op_package打包算子并分别add_subdirectory(framework)、op_host、op_kernel。七、算子实现Host 侧算子定义与 Tiling7.1 算子定义OpDef在 op_host/sobel_custom.cpp 中SobelCustom类继承OpDef完成算子注册输入x必选参数数据类型ge::DT_UINT8格式FORMAT_ND输出y必选参数数据类型ge::DT_UINT8格式FORMAT_ND绑定InferShape、InferDataType回调输出数据类型与输入一致AICore 配置开启动态编译、动态 Format、动态 Rank、动态 Shape 支持通过AddConfig(kirinx90, aicore_config)注册 Kirin X90 处理器配置。7.2 Tiling 实现Tiling 策略在 sobel_custom.cpp 的 TilingFunc 中计算核心参数如下常量值含义USE_CORE_NUM1使用 1 个 AI CoreBUFFER_NUM2双缓冲队列深度tileNum8循环分块数量h9Tile 高含 3×3 卷积所需的 halo 重叠行w256Tile 宽c3通道数Tiling 计算逻辑由输入GetShapeSize()得到totalLengthblockLength totalLength / USE_CORE_NUMtileLength h * w * cgrayLength h * w并读取输入 Shape 的 H、W 维度写入 Tiling 数据最后context-SetBlockDim(1)设定单核执行Workspace 大小为 0。7.3 TilingData 结构sobel_custom_tiling.h 通过BEGIN_TILING_DATA_DEF / END_TILING_DATA_DEF宏定义并在REGISTER_TILING_DATA_CLASS(SobelCustom, SobelCustomTilingData)注册TILING_DATA_FIELD_DEF(uint32_t, size); TILING_DATA_FIELD_DEF(uint32_t, totalLength); TILING_DATA_FIELD_DEF(uint32_t, blockLength); TILING_DATA_FIELD_DEF(uint32_t, tileLength); TILING_DATA_FIELD_DEF(uint32_t, grayLength); TILING_DATA_FIELD_DEF(uint32_t, tileNum); TILING_DATA_FIELD_DEF(uint32_t, H); TILING_DATA_FIELD_DEF(uint32_t, W);这些字段在 Kernel 侧通过GET_TILING_DATA取出用于驱动核函数的搬运与计算。八、算子实现Kernel 侧CopyIn / Compute / CopyOut核函数主体位于 op_kernel/sobel_custom.cpp模板类KernelSobelCustomuint8_t采用 Ascend C 经典的搬入-计算-搬出三段流水结构配合双缓冲队列TQueQuePosition::VECIN, BUFFER_NUM隐藏搬运与计算延迟。8.1 数据分块与搬入CopyInProcess()中按cntH CeilDiv(H, h)、cntW CeilDiv(W, w)将输入切分为h × w的 Tile逐块执行CopyIn → Compute → CopyOut。CopyIn针对不同位置首列、中间块、末列、最后一行分别构造DataCopyExtParams通过srcStride/dstStride处理行间跳步并使用DataCopyPad统一搬运带 halo 重叠3×3 卷积需要相邻 Tile 共享边界行/列的数据块。CeilDiv与Min工具函数定义在 sobel_custom_base.h。8.2 计算流程ComputeCompute()是算法核心代码 中的计算链如下NHWC → NCHW 转置AscendC::TransposeTRANSPOSE_NHWC2NCHW将输入转换为通道分离布局便于逐通道处理类型提升AscendC::Cast将 uint8 转为 halfFP16提升中间计算精度RGB2Gray利用Muls按0.299 / 0.587 / 0.114加权三个通道后Add合并得到灰度图Sobel-x 计算CreateVecIndex生成索引序列[2, 3, ..., w-1]Muls缩放后Cast为 uint32由于data[i][j1]、data[i][j2]等偏移地址不满足 32 字节对齐代码使用AscendC::Gather按索引取数再以Muls、Add完成-1·P0 - 2·P1 - 1·P2 2·P4 1·P5 2·P6 1·P7的卷积累加Sobel-y 计算同理生成两套索引[1..w-2]与[2..w-1]通过Gather处理非对齐访问完成垂直方向梯度梯度融合Abs(dx)、Abs(dy)取绝对值后Add得到|dx| |dy|量化输出Cast(..., CAST_CEIL)将 half 结果向上取整转回 uint8。代码注释明确说明了使用Gather的原因由于data[i][j 2]的地址不是 32Byte 对齐的所以需要用 gather这正是向量指令对地址对齐要求的典型处理手法。8.3 结果搬出CopyOutCopyOut按h-2 × w-2的有效输出块剔除 halo 行/列组织DataCopyExtParams对不同边界块左上、最后一列、最后一行分别设置blockCount、blockLen与srcStride/dstStride最终DataCopyPad写回 Global Memory。8.4 核函数入口extern C __global__ __aicore__ void sobel_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tiling_data, tiling); KernelSobelCustomuint8_t op; op.Init(x, y, tiling_data); op.Process(); }九、ONNX 框架插件适配为使 ATC 工具能够将 ONNX 模型中的SobelCustom节点映射到本自定义算子需要在 framework/onnx_plugin/sobel_custom_plugin.cc 中完成 GE 侧注册REGISTER_CUSTOM_OP(SobelCustom) // 算子注册名 .FrameworkType(ONNX) // 原始框架类型 .OriginOpType(ai.onnx::11::SobelCustom) // 原始框架中的算子类型 .ParseParamsByOperatorFn(ParseParamSobelCustom);ParseParamSobelCustom当前直接返回SUCCESS该算子无额外属性需要解析。对应的单算子 ONNX 模型由 create_onnx.py 生成使用 ONNX Helper 构造输入[1, 763, 1024, 3]UINT8、输出[1, 1, 761, 1022]UINT8的SobelCustom节点opset 版本 11输出SobelCustom.onnx供 ATC 转换。十、编译安装与调试验证10.1 环境准备编译前请参考环境准备文档完成环境搭建确保开发套件 Ascend-cann-toolkit 安装完成并确认算子目录 CMakePresets.json 中ASCEND_CANN_PACKAGE_PATH字段设置为正确的 toolkit 安装路径一般为${install_path}/cann。10.2 编译安装在算子工程目录ops/ascendc/src/sobel_custom下执行# 为脚本添加执行权限 chmod x build_and_install.sh # 执行编译安装 ./build_and_install.shbuild_and_install.sh 的主要流程若未设置ASCEND_HOME_PATH默认取/usr/local/Ascend/cann并source其setenv.bash导出ASCEND_TENSOR_COMPILER_INCLUDE与LD_PRELOADlibmmpa.so按实际路径修正CMakePresets.json中的 CANN 路径以cmake --presetdefault配置随后依次构建binary与package目标在build_out下执行生成的custom_opp_${OS_ID}_${arch}.run --quiet完成算子包安装。更多编译细节可参考算子工程编译安装指南。10.3 单算子验证golden 数据test/gen_data.py 生成输入与参考输出固定随机种子np.random.seed(42)生成(1, 763, 1024, 3)的 uint8 随机图写入x.bin用 NumPy 按0.299r 0.587g 0.114b计算灰度图用 OpenCVcv2.Sobel(..., ksize3)分别计算 x、y 方向梯度取绝对值后求和、裁剪到[0, 255]、np.ceil转 uint8写入y.bin作为 golden 数据。该脚本同时扮演可执行规格说明的角色与 Kernel 侧 Compute() 的算子语义一一对应可用于校验 NPU 输出精度。10.4 调试与端到端转换执行test目录下的create_onnx.py与gen_data.py生成单算子 ONNX 模型、输入与 golden 数据后可通过 ATC 工具指南 将模型转换为鸿蒙端可加载的 omc 模型再调用鸿蒙维测接口完成单算子的性能与精度验证。十一、鸿蒙端端到端部署Soble 应用11.1 应用工程模型端到端部署参考 Soble 鸿蒙应用说明。应用工程位于 harmony_infer/harmony_os_next/Soble核心代码entry/src/main/ ├── cpp/ │ ├── SobelCustom.cpp // native 层接口实现 │ ├── HIAIModelManager.cpp // 模型管理类实现 │ ├── HIAIModelManager.h // 模型管理类定义 │ ├── CMakeLists.txt // native 层编译配置 │ └── types/libentry/ // native 接口注册 ├── ets/ │ ├── entryability/EntryAbility.ets │ └── pages/Index.ets // 主界面CPU/NPU 推理切换 └── resources/ ├── base/media/ // cup.jpg、guitar.jpg 等测试图片 └── rawfile/SobelCustom.omc // ATC 转换后的模型文件11.2 部署步骤使用 DevEco 构建应用前将 ATC 转换后的SobelCustom.omc模型文件放置到应用entry/src/main/resources/rawfile目录下再进行应用构建和安装在手机主屏幕点击NPU 推理应用自动加载模型并执行推理展示处理后的图片与模型处理时间运行截图见文首NPU 处理耗时以界面实测显示为准点击CPU 推理可将图片交由 CPU 处理便于对比 CPU 与 NPU 的处理时间点击Click for next image切换下一张测试图片退出应用时自动卸载模型。11.3 底层 API 调用链应用依赖 CANN 动态库libhiai_foundation.so与 AI 领域公共动态库libneural_network_core.sonative 侧通过 HiAI Foundation 与 NNCore 定义的一组 API 完成模型加载 → 编译 → 执行模型兼容性检查与加载HMS_HiAICompatibility_CheckFromBuffer、OH_NNCompilation_ConstructWithOfflineModelBuffer编译与设备设置OH_NNCompilation_SetDevice、HMS_HiAIOptions_SetBandMode、HMS_HiAIOptions_SetModelDeviceOrder、OH_NNCompilation_BuildExecutor 构造与张量管理OH_NNExecutor_Construct、OH_NNExecutor_CreateInputTensorDesc/CreateOutputTensorDesc、OH_NNTensor_Create、OH_NNTensor_GetDataBuffer/GetSize同步推理OH_NNExecutor_RunSync(executor, inputTensor[], inputCount, outputTensor[], outputCount)资源释放OH_NNTensor_Destroy、OH_NNTensorDesc_Destroy、OH_NNCompilation_Destroy、OH_NNExecutor_Destroy。OH_NNExecutor_RunSync是推理的关键入口将待处理的图片数据填充到输入 Tensor 后同步执行输出 Tensor 即 Sobel 边缘检测结果。该应用的模型管理实现可参考 HIAIModelManager.cpp界面逻辑参考 Index.ets。11.4 约束与限制本示例仅支持标准系统上运行支持华为手机、平板和 2in1 设备HarmonyOS 系统HarmonyOS 5.0.1 Release 及以上DevEco Studio 版本DevEco Studio 6.0.0 Release 及以上HarmonyOS SDK 版本HarmonyOS 6.0.0 Release SDK 及以上。十二、总结本样例完整展示了从 Ascend C 算子实现到鸿蒙端侧推理的整条链路算法侧Sobel 算子的RGB2Gray → Sobel-x/y → |dx||dy|全流程向量化通过Transpose、Cast、Gather、Muls、Add、Abs等 Ascend C 指令实现并正确处理了非 32 字节对齐访问工程侧标准三段式算子工程framework / op_host / op_kernel配合build_and_install.sh一键编译安装ONNX 插件注册打通 ATC 转换验证侧create_onnx.pygen_data.py提供可复现的单算子模型与 golden 数据部署侧ATC 产物接入鸿蒙 Soble 应用通过 HiAI Foundation / NNCore API 在 NPU 上完成推理实现 CV 前处理从 CPU 到 NPU 的迁移。对于图像类端侧推理场景SobelCustom 提供了一种可复用的范式将灵活多变的 CV 前处理以自定义算子形式固化到 NPU 执行流水从而减少 CPU-NPU 数据往返、提升端侧推理整体效率。相关文档与源码索引本文关联文档custom-npu_sobel.md算子工程ops/ascendc/src/sobel_custom编译安装指南ops/ascendc/src/README.md环境准备docs/quick_install.mdATC 工具docs/atc_tools_guide.md鸿蒙端部署Soble/readme_cn.md算子开发通用指南ascendc_develop_guide.md【免费下载链接】cann-recipes-harmony-infer本项目为鸿蒙开发者提供基于CANN平台的业务实践案例方便开发者参考实现端云能力迁移及端侧推理部署。项目地址: https://gitcode.com/cann/cann-recipes-harmony-infer创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考