企业官网建设流程全解析

最近在昇腾设备上做模型训练，发现框架自带算子覆盖不了我需要的奇怪融合逻辑，就开始倒腾MindSpore的Custom原语。折腾了一圈下来，最大的感触是：自定义算子的前向逻辑并不难，难的是把func_type配对、把bprop写对，让反向传播能顺滑地跑起来。这篇东西想把这段时间的实践经验整理出来，围绕“用Custom原语调用Ascend C算子”这个场景，重点聊聊func_type和bprop的配置思路，以及我在实际调试中踩过的坑。

我自己是在做某个跨平台系统的时候遇到这个需求的。当时实验模型里需要把乘法和缩放在一个kernel里融合掉，降低访存开销。虽然MindSpore自带的Mul、Muls也能做，但连续调用多个算子会有多次host-device通信，性能并不理想。后来决定用Ascend C写一个自定义kernel，再通过Custom原语把它接入训练图。结果发现，如何在Python侧正确声明这个算子的执行方式与梯度逻辑，才是真正让人头疼的地方。

如果你已经在用MindSpore，并且打算在昇腾设备上写自己的融合算子，这篇文章应该能帮你省下不少排查时间。我会从一个最简单的my_mul算子说起，逐步拆解func_type的选型逻辑、bprop的编写规则，最后给出一套可以直接跑通前向和反向的完整示例。

1. 先搞清楚Custom原语到底在做什么

1.1 从一次自定义算子需求说起

MindSpore的Custom原语，本质上是一个“外部算子接入桥”。它允许你把一个用Ascend C、C++甚至Python写的函数，包装成MindSpore计算图里的一个普通算子节点。框架不关心你这个函数内部怎么实现，只在乎几个声明信息：函数体是什么、输入输出长什么样、梯度怎么算。

我最初的理解很朴素：把Ascend C编译好的so文件加载进来，然后像调用普通算子一样调用它。但实际接入后才发现，Custom原语对“如何执行”这件事分得非常细，func_type就是干这个的。它不只是一个格式枚举，而是决定了这个算子在哪类硬件上执行、以什么方式被加载、能否参与后续的编译优化。

举个例子，如果func_type配成pyfunc，那MindSpore会把你的Python函数当作一个普通Python op来执行。这种方式在CPU调试阶段很灵活，但一旦进入昇腾训练，性能基本没法看。因为pyfunc没有算子级编译优化，也走不了图编译的融合通道，本质上是把张量数据拿到host侧用Python逻辑算一遍再搬回device，这对训练场景简直是灾难。

而我用Ascend C写的算子，目标是直接在昇腾AI处理器上跑。所以func_type必须要能指向一个编译产物，而不是一个Python函数。这就要说到aot和mindir这两个更专业的选项了。

1.2 func_type和bprop在Custom中的角色

我觉得可以用一个生活化类比来解释func_type和bprop的关系。

func_type解决的是“MindSpore怎么找到你的算子实现，并让它在正确的设备上跑起来”。而bprop解决的是“梯度从输出传回输入时，这条路径该怎么算”。前者像外卖平台怎么找到餐厅并配送，后者像你吃完后怎么给餐厅打分评价——两者都重要，但完全独立。

在Custom算子注册时，func_type决定前向执行方式，bprop则是反向传播的钩子函数。只要你在注册时传递了bprop，MindSpore的自动微分引擎在计算梯度时就会调用它，而不会尝试对前向函数做符号求导。

这里有一个很关键的点：很多初学者以为只要写了Ascend C算子，反向传播就自动能跑。实际上不是。MindSpore不知道你的Ascend C kernel内部是什么数学公式，也就无法自动推导梯度。你必须显式地在bprop里告诉框架：给定输出梯度dout，每个输入各应该回传多少个梯度。

换句话说，前向逻辑是“黑盒”的，反向逻辑必须是“白盒”的。这个设计虽然给开发者增加了工作量，但也给了极大的自由度——你可以在bprop里实现任何梯度逻辑，甚至可以用另一个Ascend C算子来加速梯度计算。

1.3 几种常见的func_type选型对比

在实操中，我接触到的func_type主要有三种选择：pyfunc、aot、mindir。我整理了一个对比表，可以参考：

对于Ascend C编写的算子，我最终选择的是aot。原因很简单：Ascend C编译器会生成一个包含kernel入口的.so文件，aot模式可以直接加载这个.so，并在昇腾设备上执行。这个路径最短，也最贴近算子开发的本质。

mindir模式我也试过。它的思路是把自定义算子先封装成一个MindIR模型文件，再通过Custom加载。优点是MindIR带有完整的图结构信息，框架可以做更多编译期优化；但缺点也很明显：每次修改算子都要重新导出MindIR，调试周期长，而且对于只想把一个kernel塞进训练图的人来说，这明显“杀鸡用了牛刀”。

2. func_type配置：决定算子以哪种方式被调用

2.1 func_type="aot"与Ascend C编译器

在aot模式下，MindSpore会按照你给出的func字符串去对应的共享库中查找同名符号，并把它作为算子入口。这个符号必须和Ascend C代码中extern "C" __global__ __aicore__导出的函数名严格一致。

我之前第一次编译时，把函数名写成了my_mul_kernel，Python侧注册的func="my_mul"，结果一运行就报“undefined symbol”。排查了半天才发现，MindSpore按my_mul去so里找符号，但so里只有my_mul_kernel，两边对不上。所以请记住：Python侧的func字符串、Ascend C代码里的kernel函数名、以及编译产物的导出符号名，三者必须保持一致。

在开发阶段，我建议把编译好的so文件直接放在当前脚本目录下。MindSpore在aot模式下会从进程关联的自定义算子路径中搜索so文件，当前目录是其中最省事的选项。如果你把so放在其他位置，则需要通过环境变量指定搜索路径，否则即使符号名正确也会报“can not find lib”之类的错误。

另外，aot模式下的函数执行是在昇腾设备上的。这意味着你的Ascend C代码里必须正确管理Global Memory、Local Memory和搬运逻辑。MindSpore只是负责把输入张量地址传给kernel入口，剩下的都交给kernel自己处理。这也是Ascend C编程和普通C++编程最大的不同。

2.2 func_type="mindir"的适用场景

mindir模式在一些场景下依然有价值。

比如，你希望使用MindSpore的图编译优化来融合你的自定义算子，或者你的算子已经通过MindIR文件分发到不同后端（昇腾、GPU），那么mindir就是更合适的方案。因为它携带了完整的输入输出shape、dtype和计算图结构，MindSpore可以在编译期做更多推断和优化。

但对我这种“想在训练脚本里快速接入一个手工kernel”的需求来说，mindir显得太重了。每次修改kernel逻辑后，都要重新走一遍算子工程生成、编译、导出MindIR的流程，调试效率明显不如aot。所以我最终的方案是：开发调试用aot，等算子逻辑稳定之后，如果团队需要把它集成到更大的推理/训练pipeline里，再考虑封装成MindIR。

2.3 配置func_type时经常踩的坑

配置func_type看起来就是填一个字符串，实际操作中却有不少坑。我把踩过的和帮别人处理过的问题列一下，希望能帮你避开：

• 在CPU环境尝试跑Ascend C算子：aot加载的so依赖昇腾runtime，如果当前进程没有昇腾设备上下文，即使so加载成功，运行也会崩溃。先用npu-smi info之类的命令确认设备可见。
• 符号名不一致：前面提到的my_mul_kernelvsmy_mul问题，非常典型。建议编译后用nm -D libxxx.so | grep 函数名检查导出符号。
• so文件搜索路径不对：MindSpore不会去任意目录找so。最稳妥的做法是：开发阶段把so放在当前目录，部署阶段设置好自定义算子目录的环境变量。
• 编译选项与运行时版本不匹配：Ascend C编译产物的runtime版本不能高于当前CANN环境版本，否则加载时会报版本不兼容。这种问题通常在一台机器编译、另一台机器运行时暴露。
• 多个so文件符号冲突：如果同一个进程里存在两个so都导出了同名kernel函数，加载顺序会影响实际调用的实现。建议给kernel命名尽量加上项目前缀，避免冲突。

这些坑里面，我觉得“符号名不一致”是最容易踩的，也是排查成本最高的。因为错误信息有时只告诉你“undefined symbol”，却不会提示你它到底去哪个so里找。最好在编写CMake和Python注册脚本时，就把函数名统一维护在一个变量里，从源头避免不一致。

3. bprop反向传播：给Custom算子补上梯度逻辑

3.1 bprop函数的参数和返回格式

bprop是Custom原语的可选参数，它的函数签名遵循一个固定模式。我习惯这样定义：

def my_mul_bprop(x, y, out, dout): # x, y: 前向输入 # out: 前向输出 # dout: 输出梯度 dx = dout * y dy = dout * x return dx, dy

这里的参数顺序可以理解为：所有前向输入参数，接着是前向输出，最后是输出梯度。返回的是一个元组，每个元素对应一个输入张量的梯度。

有一点容易被忽略：如果前向算子有多个输出，那么dout就不是一个张量，而是一个元组。每个输出梯度都要在bprop里逐一使用。同样的，如果算子有两个以上的输入，返回的元组长度必须和输入数量一致。

如果某个输入不需要梯度，我建议不要偷懒返回None，更稳妥的做法是返回一个形状相同的零张量，必要时再用ops.stop_gradient显式截断。因为在图编译阶段，返回None有时候会触发类型推断错误，返回零张量则语义更清晰。

3.2 梯度公式推导与代码对照

对于my_mul这样的简单算子，梯度推导非常直接。

前向计算是：

• out = x * y

那么对x的偏导是y，对y的偏导是x。再乘上输出梯度dout，就得到输入的梯度：

• dx = dout * y
• dy = dout * x

把公式翻译成MindSpore代码，就是刚才展示的my_mul_bprop。

这里我想提出一个实践建议：**写bprop之前，先把梯度公式在纸上推一遍，或者用MindSpore自带算子做一个小规模的数值梯度对比验证。**不要等接入训练脚本之后才测试梯度。我曾经在一个类似out = x * scale + bias的算子中，把scale的梯度方向写反了，导致训练loss初期下降后又迅速发散，排查了很久才发现是梯度符号错误。如果当时先用数值梯度校验一遍，五分钟就能发现。

3.3 多个输入/输出场景下的bprop写法

多输入的场景刚才已经提到了。多输出则稍微复杂一点。

假设你的自定义算子输出两个结果out1, out2，那么dout会是(grad_out1, grad_out2)。你需要分别处理每个输出的梯度，最后合并返回输入梯度。

def multi_out_bprop(x, out1, out2, dout1, dout2): dx = dout1 * 2.0 + dout2 * x return (dx,)

这里dout1和dout2分别是两个输出节点的梯度。组合方式取决于前向公式。如果前向是：

• out1 = 2.0 * x
• out2 = x * x

那么：

• dx = dout1 * 2.0 + dout2 * 2.0 * x

这就是链式法则的展开。把推导写在纸上，再翻译成代码，会清晰很多。

如果遇到某些输出节点在后向传播中根本不需要梯度，你依然要在dout对应的位置上接收一个零张量，而不是跳过它，因为Python的元组索引位置不能变。

3.4 bprop与AutoDiff、stop_gradient的关系

MindSpore的自动微分在训练图构建阶段，会对每个算子查询它的bprop。如果Custom算子在注册时提供了bprop，自动微分就会直接使用它；如果没提供，框架会尝试对该算子的前向实现做自动展开——但这对于黑盒的Ascend C代码几乎不可能成功，所以最终会报错。

这就意味着，你可以在bprop里调用MindSpore的任何算子组合，甚至可以再次调用另一个Custom算子。我实际做过一个融合算子的反向：前向是一个Ascend C kernel，反向里又调用了另外一个Ascend C kernel来加速梯度计算。bprop的存在让这种“前向和反向各自黑盒、框架只负责调度”的架构变得非常顺滑。

stop_gradient则是一个特殊的梯度阻断算子。如果某个输入的梯度不需要回传，你可以在bprop中对该梯度调用ops.stop_gradient。这样做的好处是，后续反向图构建时，这个分支会被彻底剪掉，能省掉不少计算。但请注意，stop_gradient不要用在所有输入的梯度上，否则反向就完全没有意义了。

4. 完整实操：把一个Ascend C算子接进MindSpore并跑通反向

4.1 编写Ascend C算子

下面以my_mul为例，给出一个简化但结构完整的Ascend C算子实现。这个例子计算out = x * y，输入和输出都是半精度浮点张量。

#include "kernel_operator.h" using namespace AscendC; constexpr int32_t BUFFER_NUM = 2; class KernelMyMul { public: __aicore__ inline KernelMyMul() {} __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR out, uint32_t totalLength, uint32_t blockLength) { this->totalLength = totalLength; this->blockLength = blockLength; xGm.SetGlobalBuffer((__gm__ half*)x, totalLength); yGm.SetGlobalBuffer((__gm__ half*)y, totalLength); outGm.SetGlobalBuffer((__gm__ half*)out, totalLength); pipe.InitBuffer(inQueueX, BUFFER_NUM, blockLength * sizeof(half)); pipe.InitBuffer(inQueueY, BUFFER_NUM, blockLength * sizeof(half)); pipe.InitBuffer(outQueue, BUFFER_NUM, blockLength * sizeof(half)); } __aicore__ inline void Process() { uint32_t blockIndex = GetBlockIdx(); uint32_t blockOffset = blockIndex * blockLength; LocalTensor<half> xLocal = inQueueX.AllocTensor<half>(); LocalTensor<half> yLocal = inQueueY.AllocTensor<half>(); LocalTensor<half> outLocal = outQueue.AllocTensor<half>(); DataCopy(xLocal, xGm[blockOffset], blockLength); DataCopy(yLocal, yGm[blockOffset], blockLength); for (uint32_t i = 0; i < blockLength; i++) { outLocal.SetValue(i, xLocal.GetValue(i) * yLocal.GetValue(i)); } DataCopy(outGm[blockOffset], outLocal, blockLength); inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal); outQueue.EnQue(outLocal); } private: GlobalTensor<half> xGm, yGm, outGm; TPipe pipe; TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX; TQue<QuePosition::VECIN, BUFFER_NUM> inQueueY; TQue<QuePosition::VECOUT, BUFFER_NUM> outQueue; uint32_t totalLength; uint32_t blockLength; }; extern "C" __global__ __aicore__ void my_mul(GM_ADDR x, GM_ADDR y, GM_ADDR out) { uint32_t totalLength = GetTensorNum(); uint32_t blockLength = GetBlockSize(totalLength); KernelMyMul op; op.Init(x, y, out, totalLength, blockLength); op.Process(); }

这段代码里，GetTensorNum和GetBlockSize是Ascend C框架提供的辅助函数，用于在核函数入口处计算张量长度和每个核负责的分块长度。实际工程中你需要从算子描述信息中读取shape，我这里做了简化处理。

注意，这里的实现只是示意，没有处理尾部剩余数据、多核分块的完整逻辑。真正上生产前，需要补全边界分支。但结构上已经足够说明一个Ascend C kernel的骨架。

4.2 编译算子并准备.so文件

编译Ascend C算子需要CANN提供的工具链。我习惯用CMake组织算子工程，关键的CMake配置大致长这样：

cmake_minimum_required(VERSION 3.16) project(my_mul CXX) set(CMAKE_CXX_STANDARD 17) set(CMAKE_CXX_FLAGS "-fPIC -O2") include_directories(${ASCEND_HOME}/include) add_library(my_mul SHARED my_mul.cpp)

编译命令很简单：

mkdir -p build && cd build cmake .. && make -j8

编译完成后，当前目录会生成libmy_mul.so。用下面的命令确认导出符号：

nm -D libmy_mul.so | grep my_mul

只要能看到my_mul符号，就说明命名没问题。开发阶段直接把so放在Python脚本目录下，后续可以直接加载。

4.3 Python侧Custom注册与bprop实现

Python侧的核心代码如下：

import numpy as np import mindspore as ms import mindspore.ops as ops import mindspore.common.dtype as mstype from mindspore.ops import Custom def my_mul_out_shape(x, y): return x.shape def my_mul_out_dtype(x, y): return x.dtype def my_mul_bprop(x, y, out, dout): dx = dout * y dy = dout * x return dx, dy my_mul_op = Custom( func="my_mul", out_shape=my_mul_out_shape, out_dtype=my_mul_out_dtype, func_type="aot", bprop=my_mul_bprop )

这段代码里，func字符串必须与Ascend C代码导出的符号名一致。out_shape和out_dtype接收的函数，用来告诉MindSpore这个算子的输出形状和数据类型，它们会被用于图编译期的shape推断。

bprop就是反向传播钩子。my_mul_bprop接收前向输入x, y、前向输出out和输出梯度dout，返回两个梯度。梯度返回顺序必须与输入参数顺序一致，也就是(dx, dy)。

4.4 前向+反向端到端验证

算子注册完成后，可以写一个简单的验证脚本，分别测试前向和反向是否正常：

import mindspore as ms import mindspore.ops as ops import numpy as np x = ms.Tensor(np.array([[1.0, 2.0], [3.0, 4.0]]), ms.float16) y = ms.Tensor(np.array([[5.0, 6.0], [7.0, 8.0]]), ms.float16) # 前向验证 out = my_mul_op(x, y) print(out) # 反向验证 grad_fn = ms.grad(my_mul_op, grad_position=(0, 1)) dx, dy = grad_fn(x, y) print(dx) print(dy)

ms.grad会利用my_mul_bprop计算梯度。如果一切正常，dx应该等于dout * y，dy应该等于dout * x。在初始验证时，dout默认是全1张量，所以可以直观对比结果。

这里有一个小技巧：先跑前向，确认so能被正确加载、结果数值正确，再去验证反向。不要一上来就同时验证两个方向，否则出错了很难分清是前向加载的问题，还是反向梯度的问题。

5. 常见问题与排查技巧实录

5.1 func_type配置不当导致算子无法加载

常见报错包括：

• RuntimeError: func my_mul is not registered
• RuntimeError: can not find symbol my_mul
• RuntimeError: can not find lib

排查思路按顺序来：

1. 确认so文件存在且在Python进程的搜索路径中。
2. 用nm -D libmy_mul.so | grep my_mul确认符号名。
3. 确认编译产物与当前CANN运行环境匹配。
4. 确认进程能访问昇腾设备，简单方法是在脚本开头执行一个普通的MindSpore算子，看是否正常。

有一次我遇到一个奇怪情况：so文件路径正确、符号也存在，但依然报找不到。后来发现是环境变量里配置了多个自定义算子目录，MindSpore在某个优先目录里先找到了另一个同名so，加载了错误的库。解决方法是精简目录配置，只保留一个可信路径。

5.2 bprop梯度不匹配报错

如果bprop返回的梯度数量或形状不对，MindSpore会在图编译阶段报错。常见的报错信息是：

• For 'Custom', the number of gradients returned from bprop should be N
• The shape of gradient should be same as corresponding input

第一种错误通常是因为返回的梯度数量与输入数量不一致。例如：有2个输入，但bprop只返回了一个梯度。解决办法是检查返回语句，确保返回元组的长度等于输入数量。

第二种错误则说明梯度的shape与输入shape对不上。常见于前向函数里做了reshape或broadcast，而bprop里没有对应处理。比如输入是[2, 3]，前向输出是[6]，那么bprop计算出的梯度也得是[2, 3]，不能是[6]。如果前向做了reshape，反向需要先算出梯度再reshape回输入形状。

5.3 反向不执行的几种情况

有时候算子前向正常，loss也正常下降，但某一层参数就是不更新，大概率是反向逻辑没有生效。我总结了几个典型情况：

• 没有开启训练模式：有些自定义算子的bprop只在nn.Cell的training模式下参与梯度构建。如果推理模式跑，自动微分可能跳过梯度计算。
• 输入没有requires_grad：在MindSpore中，默认对Parameter创建requires_grad，但对普通Tensor可能需要显式标记。如果输入不要求梯度，bprop即使定义正确也不会被调用。
• 在图模式与PyNative模式下行为不一致：我遇到过在图模式下bprop正常、切到PyNative模式就报错的情况。原因是PyNative模式下bprop中的某些算子组合没有被正确追踪，建议统一使用Graph模式跑训练验证。

排查反向是否执行，可以用一个笨办法：在bprop里给dx加上一个极大的常数，比如dout * y + 9999.0。如果反向被调用了，loss一定会剧烈变化；如果没有变化，说明bprop根本没被走到。定位后，再把常数删掉。

5.4 性能排查思路

如果前向和反向都正确，但训练速度比预期慢，建议从这几个方向排查：

• 确认算子是否真的在昇腾设备上执行：最简单的方法是在训练脚本里关闭pyfunc模式，改用aot。如果性能有明显提升，说明之前的算子在host侧跑，瓶颈卡在数据搬运上。
• 检查kernel的分块策略：Ascend C算子的性能很大程度上取决于每个核处理的数据量。分块太小，核间通信开销大；分块太大，可能超出局部内存容量。建议根据输入shape和昇腾硬件核数，用不同分块参数做benchmark。
• 看host-device数据拷贝次数：自定义算子如果每次调用都触发大量host-device拷贝，性能会非常差。结合昇腾profiler工具，观察算子间的数据流，优化kernel的输入输出布局。
• 避免bprop里使用过多Python级算子：虽然bprop可以用MindSpore算子自由组合，但组合过于复杂会拉低反向计算效率。如果反向计算量很大，建议再写一个Ascend C反向kernel，在bprop里通过Custom调用它。这个思路在做深度自研算子时非常实用。

结尾

折腾了这段时间，我个人最大的体会是：Custom原语把“接入一个外部算子”的门槛降得很低，但真正决定开发效率的，是对func_type和bprop这两个参数的理解深度。

func_type的配置本质上是告诉框架“你在用什么方式调用实现”，它决定了算子的执行路径和性能前景。对于Ascend C算子，aot几乎是最优解，因为它直接加载编译产物，省掉了中间转换流程。bprop则是反向传播的桥梁，前向可以黑盒，反向必须白盒。只要你能把梯度公式写清楚，MindSpore就能自动完成反向图的构建和调度。

最后分享一个小技巧：在你第一次写好bprop后，用数值微分法做一次梯度校验。具体做法就是构造随机输入，用(f(x+epsilon) - f(x-epsilon)) / (2*epsilon)近似每个输入的梯度，再和bprop返回的梯度做对比。这一步能帮你发现八成以上的梯度推导错误，成本只有几分钟。等算子逻辑稳定了，再把这个校验脚本放到自动化测试里，以后每次改kernel都能跑一遍，防止回归。

这个内容后续还可以扩展到更复杂的融合算子，比如带reduce操作或者涉及矩阵乘法的场景。那类算子的梯度推导会更依赖链式法则，但配置思路是完全一致的——希望这篇经验能帮你少走一些弯路。