从0到1掌握Ascend C算子工程开发方式

（如。

2401_89199769

2276人浏览 · 2025-10-24 16:14:21

2401_89199769 · 2025-10-24 16:14:21 发布

（一）两种算子开发流程

快速开发调试方式（Ascend C 算子）：聚焦算子核函数开发，通过Kernel 直调实现算子调用运行。流程为：环境准备→矢量 / 矩阵编程（Kernel 实现）→基于内核调用符进行算子运行验证→算子调试调优。
标准开发调试方式（Ascend C 算子）：需完成Host 侧 + Device 侧 + 应用程序全链路开发，通过单算子 API（如 ACLNN、ACLOP）或 PyTorch Adapter 调用。流程为：环境准备（CANN 软件安装）→算子分析→创建算子工程→Kernel 算子实现→Host 侧算子实现→编译部署→ST 测试→算子调用（AscendCL 单算子 / 网络中算子调用）。
模式对比：\
（二）算子工程之 “算子分析”

需明确算子的数学表达式、计算逻辑，以及输入输出张量的shape、dtype、format等参数（如示例中算子输入 A、B 和输出 C 的维度、数据类型定义）。计算逻辑基于 Ascend C 的张量运算，需分析 Local Memory 与 Global Memory 的调度、算子并行性及数据依赖关系。

（三）算子工程之 “创建算子工程”

通过 CANN 开发套件工具mgcGcn，基于算子原型定义（如add_custom.json）创建工程，自动生成 Host 侧代码文件、Kernel 实现文件及编译配置文件。示例命令：mgcGcn add -i add_custom.json -c core ascendclan -o out AddCustom。

（四）Host&Device 架构
Host：与 Device 连接的 X86/ARM 服务器，负责任务调度、数据准备，利用 Device 的 NN Network 模块算力完成任务。
Device：搭载昇腾硬件（如华为 AI 加速卡）的设备，通过 PCIe 接口与 Host 通信，提供 AI 算力加速能力。
（五）Host 侧算子实现与 Tiling 技术
Host 侧算子实现核心：聚焦 Tiling（数据分块）、Shape 推导（张量维度推导）、算子注册（封装算子信息），涉及tiling头文件（如xxx_custom_tiling.h）、业务代码文件（如add_custom.cpp）等。
Tiling 技术：
- 概念：因 Local Memory 空间有限，将输入数据分块处理的技术，分块过程称为 Tiling 实现。
- 实现流程：
  - Host 侧：定义 Tiling 结构体（描述分块规则）→实现 Tiling 函数（处理分块逻辑）。
  - Kernel 侧：接收 Tiling 信息，基于分块数据执行计算。
- 场景适配：
  - 固定 Shape：输入大小固定，提前确定分块数 / 次数。
  - 动态 Shape：通过输入参数动态推导 shape，适配不同输入规模。
  - （一）算子开发流程对比流程图

（二）Tiling 技术实现流程图

（三）算子工程模块拆解表

（六）固定与动态 Shape 对比及实现

KernelAdd 类对比：固定 Shape 场景下，KernelAdd类需硬编码固定输入输出维度；动态 Shape 场景则通过参数推导适配不同输入规模，代码中包含shape动态判断逻辑。
Init 函数差异：固定 Shape 的Init函数直接初始化固定维度资源；动态 Shape 的Init函数需调用getTensorShape等 API 推导shape后，再初始化对应资源。

（七）Host 侧 Shape 推导

概述：通过 Tensor 的shape和data type推导，在算子执行前明确输入输出张量的形状、数据类型及格式，保障网络推理流程的正确性。
核心作用：
1. 参数校验：提前发现输入参数不匹配，提升程序健壮性。
2. 输出张量描述：推导输出张量的形状、数据类型、格式等信息。
3. 内存优化：静态分配内存，避免动态内存分配的性能开销。

（八）算子包编译与部署

编译流程：
1. 修改算子工程的CMakeLists.txt和CMakePresets.json，配置ASCEND_CANN_PACKAGE_PATH等参数。
2. 执行工程下的build.sh脚本，编译完成后在build_out目录生成算子文件（如custom_opp.so）。
部署步骤：在算子包路径下执行命令./custom_opp_<target_os>_<target_architecture>.run，将算子包安装到环境 OP 库目录（如/usr/local/Ascend/ascend-toolkit/latest/opp/endor/customize）。

（九）Ascend C 算子中级认证考点（Tiling 与算子类）

Tiling 结构体分析：需掌握 Tiling 结构体定义（如SigmoidCustomTiling），包含分块维度（TILE_DIM）、分块数量（TILE_NUM）等参数，通过getTilingData传递给 Kernel 侧。
Kernel 侧 Tiling 使用：Kernel 函数需接收 Tiling 信息作为参数（如__global__ void add_custom(..., __gm__ Tiling tiling)），并基于分块数据执行计算（如通过tiling.getData<TILING_DATA_FIELD_OFF>(idx)获取分块数据）。
KernelSigmoid 类实现：需掌握算子类结构，包含Process、Compute等方法，实现 Sigmoid 算子的计算逻辑，同时适配 Tiling 分块逻辑。
固定与动态 Shape 实现差异对比表
Host 侧 Shape 推导流程图