训练营简介：2025年昇腾CANN训练营第二季，基于CANN开源开放全场景，推出0基础入门系列、码力全开特辑、开发者案例等专题课程，助力不同阶段开发者快速提升算子开发技能。获得Ascend C算子中级认证，即可领取精美证书，完成社区任务更有机会赢取华为手机，平板、开发板等大奖。

一、什么是算子

在深度学习模型中，算子（Operator） 是构成神经网络层或节点的基本计算逻辑。例如，卷积（Convolution）、批归一化（BatchNorm）、激活函数（ReLU）等都是常见的算子。不同的框架（如Caffe、TensorFlow）对算子的实现和组合方式有所不同，但本质都是对输入张量（Tensor）进行数学运算，输出新的张量。

二、为什么需要自定义算子？—— Ascend C的用武之地（优势）

虽然主流框架提供了丰富的预置算子，但在以下场景中，我们仍需自定义算子：

框架不支持：转换模型时遇到不支持的算子；

性能优化：替换性能低的算子，提升推理/训练速度；

算法定制：实现特定数学运算（如最大值查找、数据类型转换）；

模型调优：针对硬件特性重新设计算子逻辑。

三、核函数

核函数是设备侧（AI Core）执行的入口函数，语法如下：

__global__ __aicore__ void kernel_name(GM_ADDR x, GM_ADDR y, GM_ADDR z);

__global__ 表示该函数为核函数；

__aicore__ 表示在AI Core上执行；

参数通常为指向Global Memory的指针。

函数调用方式:

kernel name<<<blockDim,l2ctrl,stream>>>(argument list);

blockDim，规定了核函数将会在几个核上执行，每个执行该核函数的核会被分配一个逻辑ID，表现为内置变量block idx，编号从0开始，可为不同的逻辑核定义不同的行为，可以在算子实现中使用

12ctr，保留参数，暂时设置为固定值nullptr

stream，类型为acitStream,stream是一个任务队列，应用程序通过stream来管理任务的并行

四、三级流水任务（Pipeline Stages）

Ascend C将算子执行流程划分为三个并行任务：

阶段                 任务                           说明

Stage1             CopyIn                      将数据从Global Memory搬入Local Memory

Stage2             Compute                   在Local Memory中执行计算（Vector/Cube指令）

Stage3             CopyOut                    将结果从Local Memory搬回Global Memory

通过双缓冲（Double Buffer） 技术，实现数据搬运与计算的并行，显著提升效率。

五、算子的执行模式

https://i-blog.csdnimg.cn/blog_migrate/041fb07101d8383b8222ccccccdb6d17.png

支持两种模式运行：

CPU模式：用于调试与验证逻辑；

NPU模式：在昇腾硬件上实际执行。

六、总结

Ascend C通过其分层API、流水并行模型和孪生调试机制，极大地降低了算子开发的门槛。无论是初学者还是有经验的AI开发者，都能快速上手并实现高性能的自定义算子。