昇腾 Ascend C 算子开发：从语法基础到全链路编程实践

高层 API：算法高效封装

高层 API 封装了 Matmul、Softmax 等常用算法逻辑。借助'对象化封装 + 流程化调用'模式，将复杂算法转化为简洁的 API 调用，既减少重复开发工作，又大幅提升效率。

算子开发实战流程

核函数定义与算子类实现

使用 __global__ __aicore__ 限定符定义设备侧入口，调用算子类的初始化与处理函数。算子类通过 CopyIn、Compute、CopyOut 三个流水任务实现逻辑，结合 Queue 做任务同步、Pipe 做内存管理。

Init 方法实现

KernelAdd 类的 Init 方法主要完成两件事：

1. 通过 block_idx 给当前 AI Core 划分全局内存数据，把 xGM、yGM、zGM 绑定到当前核心的处理区域，实现多核并行。
2. 用 pipe 初始化输入输出队列的双缓冲内存，按 TILE_LENGTH 分配内存块，队列深度设为 BUFFER_NUM，支撑流水线并行。

Process 方法实现

Process() 方法通过循环调度 CopyIn、Compute、CopyOut 三个阶段，结合双缓冲实现流水线并行：

• CopyIn：将全局内存数据搬入本地并通过队列同步。
• Compute：执行向量加法并将结果入队。
• CopyOut：将结果搬回全局内存。

借助双缓冲和流水线机制，让数据搬运与计算并行，有效提升昇腾 AI Core 的 Vector 单元利用率。

主机侧逻辑与数据生成

主机侧 main.cpp

包含两种模式的逻辑：

• CPU 模式：用于算法逻辑快速验证，不依赖昇腾硬件，直接在 CPU 上执行计算。
• NPU 模式：用于在昇腾 NPU 硬件上执行算子，充分利用硬件加速能力。

数据生成脚本 gen_data.py

基于 NumPy 编写，用于生成 AddCustom 算子的输入数据和真值数据，为功能验证提供标准化参考。

• 实现逻辑：生成两个形状为 (8, 2048) 的 float16 随机输入数据 input_x 和 input_y（数值范围 1-100），计算加法算子的真值数据 golden（即 input_x + input_y），并以二进制格式保存到指定路径。
• 用途：为算子的开发、调试和测试提供一致的输入源与真值基准，确保计算结果的准确性。

工程执行与验证

Ascend C 算子的简易工程执行方式是将编译和执行命令封装到 run.sh 脚本中，通过不同参数实现 CPU 和 NPU 模式下的调试。

• CPU 模式：执行 bash run.sh -v Ascend910x -r cpu，通过 md5 校验 output_z.bin 与 golden.bin 一致性验证执行成功。架构依赖 Host APP、CPU 调用库、算子 kernel 程序和 AscendC 类库。
• NPU 模式：执行 bash run.sh -v Ascend910x -r npu，通过 md5 校验验证编译、执行成功。架构依赖 Host APP、AscendCL API 库及设备侧算子 kernel 程序，实现主机 - 设备交互与功能验证。